//
// Copyright (c) 2017 The Khronos Group Inc.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
//    http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
//
#ifndef _COMMON_H_
#define _COMMON_H_

#include "harness/testHarness.h"
#include "harness/typeWrappers.h"
#include "harness/ThreadPool.h"

#include "host_atomics.h"

#include <vector>
#include <sstream>

#define MAX_DEVICE_THREADS (gHost ? 0U : gMaxDeviceThreads)
#define MAX_HOST_THREADS GetThreadCount()

#define EXECUTE_TEST(error, test)\
  error |= test;\
  if(error && !gContinueOnError)\
  return error;

enum TExplicitAtomicType
{
    TYPE_ATOMIC_INT,
    TYPE_ATOMIC_UINT,
    TYPE_ATOMIC_LONG,
    TYPE_ATOMIC_ULONG,
    TYPE_ATOMIC_FLOAT,
    TYPE_ATOMIC_DOUBLE,
    TYPE_ATOMIC_INTPTR_T,
    TYPE_ATOMIC_UINTPTR_T,
    TYPE_ATOMIC_SIZE_T,
    TYPE_ATOMIC_PTRDIFF_T,
    TYPE_ATOMIC_FLAG
};

enum TExplicitMemoryScopeType
{
    MEMORY_SCOPE_EMPTY,
    MEMORY_SCOPE_WORK_GROUP,
    MEMORY_SCOPE_DEVICE,
    MEMORY_SCOPE_ALL_DEVICES, // Alias for MEMORY_SCOPE_ALL_SVM_DEVICES
    MEMORY_SCOPE_ALL_SVM_DEVICES
};

extern bool gHost; // temporary flag for testing native host threads (test verification)
extern bool gOldAPI; // temporary flag for testing with old API (OpenCL 1.2)
extern bool gContinueOnError; // execute all cases even when errors detected
extern bool gNoGlobalVariables; // disable cases with global atomics in program scope
extern bool gNoGenericAddressSpace; // disable cases with generic address space
extern bool gUseHostPtr; // use malloc/free instead of clSVMAlloc/clSVMFree
extern bool gDebug; // print OpenCL kernel code
extern int gInternalIterations; // internal test iterations for atomic operation, sufficient to verify atomicity
extern int gMaxDeviceThreads; // maximum number of threads executed on OCL device
extern cl_device_atomic_capabilities gAtomicMemCap,
    gAtomicFenceCap; // atomic memory and fence capabilities for this device

extern const char *get_memory_order_type_name(TExplicitMemoryOrderType orderType);
extern const char *get_memory_scope_type_name(TExplicitMemoryScopeType scopeType);

extern cl_int getSupportedMemoryOrdersAndScopes(
    cl_device_id device, std::vector<TExplicitMemoryOrderType> &memoryOrders,
    std::vector<TExplicitMemoryScopeType> &memoryScopes);

class AtomicTypeInfo
{
public:
  TExplicitAtomicType _type;
  AtomicTypeInfo(TExplicitAtomicType type): _type(type) {}
  cl_uint Size(cl_device_id device);
  const char* AtomicTypeName();
  const char* RegularTypeName();
  const char* AddSubOperandTypeName();
  int IsSupported(cl_device_id device);
};

template<typename HostDataType>
class AtomicTypeExtendedInfo : public AtomicTypeInfo
{
public:
  AtomicTypeExtendedInfo(TExplicitAtomicType type) : AtomicTypeInfo(type) {}
  HostDataType MinValue();
  HostDataType MaxValue();
  HostDataType SpecialValue(cl_uchar x)
  {
    HostDataType tmp;
    cl_uchar *ptr = (cl_uchar*)&tmp;
    for(cl_uint i = 0; i < sizeof(HostDataType)/sizeof(cl_uchar); i++)
      ptr[i] = x;
    return tmp;
  }
  HostDataType SpecialValue(cl_ushort x)
  {
    HostDataType tmp;
    cl_ushort *ptr = (cl_ushort*)&tmp;
    for(cl_uint i = 0; i < sizeof(HostDataType)/sizeof(cl_ushort); i++)
      ptr[i] = x;
    return tmp;
  }
};

class CTest  {
public:
  virtual int Execute(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements) = 0;
};

template<typename HostAtomicType, typename HostDataType>
class CBasicTest : CTest
{
public:
  typedef struct {
    CBasicTest *test;
    cl_uint tid;
    cl_uint threadCount;
    volatile HostAtomicType *destMemory;
    HostDataType *oldValues;
  } THostThreadContext;
  static cl_int HostThreadFunction(cl_uint job_id, cl_uint thread_id, void *userInfo)
  {
    THostThreadContext *threadContext = ((THostThreadContext*)userInfo)+job_id;
    threadContext->test->HostFunction(threadContext->tid, threadContext->threadCount, threadContext->destMemory, threadContext->oldValues);
    return 0;
  }
  CBasicTest(TExplicitAtomicType dataType, bool useSVM) : CTest(),
    _maxDeviceThreads(MAX_DEVICE_THREADS),
    _dataType(dataType), _useSVM(useSVM), _startValue(255),
    _localMemory(false), _declaredInProgram(false),
    _usedInFunction(false), _genericAddrSpace(false),
    _oldValueCheck(true), _localRefValues(false),
    _maxGroupSize(0), _passCount(0), _iterations(gInternalIterations)
  {
  }
  virtual ~CBasicTest()
  {
    if(_passCount)
      log_info("  %u tests executed successfully for %s\n", _passCount, DataType().AtomicTypeName());
  }
  virtual cl_uint NumResults(cl_uint threadCount, cl_device_id deviceID)
  {
    return 1;
  }
  virtual cl_uint NumNonAtomicVariablesPerThread()
  {
    return 1;
  }
  virtual bool ExpectedValue(HostDataType &expected, cl_uint threadCount, HostDataType *startRefValues, cl_uint whichDestValue)
  {
    return false;
  }
  virtual bool GenerateRefs(cl_uint threadCount, HostDataType *startRefValues, MTdata d)
  {
    return false;
  }
  virtual bool VerifyRefs(bool &correct, cl_uint threadCount, HostDataType *refValues, HostAtomicType *finalValues)
  {
    return false;
  }
  virtual std::string PragmaHeader(cl_device_id deviceID);
  virtual std::string ProgramHeader(cl_uint maxNumDestItems);
  virtual std::string FunctionCode();
  virtual std::string KernelCode(cl_uint maxNumDestItems);
  virtual std::string ProgramCore() = 0;
  virtual std::string SingleTestName()
  {
    std::string testName = LocalMemory() ? "local" : "global";
    testName += " ";
    testName += DataType().AtomicTypeName();
    if(DeclaredInProgram())
    {
      testName += " declared in program";
    }
    if(DeclaredInProgram() && UsedInFunction())
      testName += ",";
    if(UsedInFunction())
    {
      testName += " used in ";
      if(GenericAddrSpace())
        testName += "generic ";
      testName += "function";
    }
    return testName;
  }
  virtual int ExecuteSingleTest(cl_device_id deviceID, cl_context context, cl_command_queue queue);
  int ExecuteForEachPointerType(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    int error = 0;
    UsedInFunction(false);
    EXECUTE_TEST(error, ExecuteSingleTest(deviceID, context, queue));
    UsedInFunction(true);
    GenericAddrSpace(false);
    EXECUTE_TEST(error, ExecuteSingleTest(deviceID, context, queue));
    GenericAddrSpace(true);
    EXECUTE_TEST(error, ExecuteSingleTest(deviceID, context, queue));
    GenericAddrSpace(false);
    return error;
  }
  int ExecuteForEachDeclarationType(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    int error = 0;
    DeclaredInProgram(false);
    EXECUTE_TEST(error, ExecuteForEachPointerType(deviceID, context, queue));
    if(!UseSVM())
    {
      DeclaredInProgram(true);
      EXECUTE_TEST(error, ExecuteForEachPointerType(deviceID, context, queue));
    }
    return error;
  }
  virtual int ExecuteForEachParameterSet(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    int error = 0;
    if(_maxDeviceThreads > 0 && !UseSVM())
    {
      LocalMemory(true);
      EXECUTE_TEST(error, ExecuteForEachDeclarationType(deviceID, context, queue));
    }
    if(_maxDeviceThreads+MaxHostThreads() > 0)
    {
      LocalMemory(false);
      EXECUTE_TEST(error, ExecuteForEachDeclarationType(deviceID, context, queue));
    }
    return error;
  }
  virtual int Execute(cl_device_id deviceID, cl_context context, cl_command_queue queue, int num_elements)
  {
    if(sizeof(HostAtomicType) != DataType().Size(deviceID))
    {
      log_info("Invalid test: Host atomic type size (%u) is different than OpenCL type size (%u)\n", (cl_uint)sizeof(HostAtomicType), DataType().Size(deviceID));
      return -1;
    }
    if(sizeof(HostAtomicType) != sizeof(HostDataType))
    {
      log_info("Invalid test: Host atomic type size (%u) is different than corresponding type size (%u)\n", (cl_uint)sizeof(HostAtomicType), (cl_uint)sizeof(HostDataType));
      return -1;
    }
    // Verify we can run first
    if(UseSVM() && !gUseHostPtr)
    {
      cl_device_svm_capabilities caps;
      cl_int error = clGetDeviceInfo(deviceID, CL_DEVICE_SVM_CAPABILITIES, sizeof(caps), &caps, 0);
      test_error(error, "clGetDeviceInfo failed");
      if((caps & CL_DEVICE_SVM_ATOMICS) == 0)
      {
        log_info("\t%s - SVM_ATOMICS not supported\n", DataType().AtomicTypeName());
        // implicit pass
        return 0;
      }
    }
    if(!DataType().IsSupported(deviceID))
    {
      log_info("\t%s not supported\n", DataType().AtomicTypeName());
      // implicit pass or host test (debug feature)
      if(UseSVM())
        return 0;
      _maxDeviceThreads = 0;
    }
    if(_maxDeviceThreads+MaxHostThreads() == 0)
      return 0;
    return ExecuteForEachParameterSet(deviceID, context, queue);
  }
  virtual void HostFunction(cl_uint tid, cl_uint threadCount, volatile HostAtomicType *destMemory, HostDataType *oldValues)
  {
    log_info("Empty thread function %u\n", (cl_uint)tid);
  }
  AtomicTypeExtendedInfo<HostDataType> DataType() const
  {
    return AtomicTypeExtendedInfo<HostDataType>(_dataType);
  }
  cl_uint _maxDeviceThreads;
  virtual cl_uint MaxHostThreads()
  {
    if(UseSVM() || gHost)
      return MAX_HOST_THREADS;
    else
      return 0;
  }

  int CheckCapabilities(TExplicitMemoryScopeType memoryScope,
                        TExplicitMemoryOrderType memoryOrder)
  {
      /*
          Differentiation between atomic fence and other atomic operations
          does not need to occur here.

          The initialisation of this test checks that the minimum required
          capabilities are supported by this device.

          The following switches allow the test to skip if optional capabilites
          are not supported by the device.
        */
      switch (memoryScope)
      {
          case MEMORY_SCOPE_EMPTY: {
              break;
          }
          case MEMORY_SCOPE_WORK_GROUP: {
              if ((gAtomicMemCap & CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP) == 0)
              {
                  return TEST_SKIPPED_ITSELF;
              }
              break;
          }
          case MEMORY_SCOPE_DEVICE: {
              if ((gAtomicMemCap & CL_DEVICE_ATOMIC_SCOPE_DEVICE) == 0)
              {
                  return TEST_SKIPPED_ITSELF;
              }
              break;
          }
          case MEMORY_SCOPE_ALL_DEVICES: // fallthough
          case MEMORY_SCOPE_ALL_SVM_DEVICES: {
              if ((gAtomicMemCap & CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES) == 0)
              {
                  return TEST_SKIPPED_ITSELF;
              }
              break;
          }
          default: {
              log_info("Invalid memory scope\n");
              break;
          }
      }

      switch (memoryOrder)
      {
          case MEMORY_ORDER_EMPTY: {
              break;
          }
          case MEMORY_ORDER_RELAXED: {
              if ((gAtomicMemCap & CL_DEVICE_ATOMIC_ORDER_RELAXED) == 0)
              {
                  return TEST_SKIPPED_ITSELF;
              }
              break;
          }
          case MEMORY_ORDER_ACQUIRE:
          case MEMORY_ORDER_RELEASE:
          case MEMORY_ORDER_ACQ_REL: {
              if ((gAtomicMemCap & CL_DEVICE_ATOMIC_ORDER_ACQ_REL) == 0)
              {
                  return TEST_SKIPPED_ITSELF;
              }
              break;
          }
          case MEMORY_ORDER_SEQ_CST: {
              if ((gAtomicMemCap & CL_DEVICE_ATOMIC_ORDER_SEQ_CST) == 0)
              {
                  return TEST_SKIPPED_ITSELF;
              }
              break;
          }
          default: {
              log_info("Invalid memory order\n");
              break;
          }
      }

      return 0;
  }
  virtual bool SVMDataBufferAllSVMConsistent() {return false;}
  bool UseSVM() {return _useSVM;}
  void StartValue(HostDataType startValue) {_startValue = startValue;}
  HostDataType StartValue() {return _startValue;}
  void LocalMemory(bool local) {_localMemory = local;}
  bool LocalMemory() {return _localMemory;}
  void DeclaredInProgram(bool declaredInProgram) {_declaredInProgram = declaredInProgram;}
  bool DeclaredInProgram() {return _declaredInProgram;}
  void UsedInFunction(bool local) {_usedInFunction = local;}
  bool UsedInFunction() {return _usedInFunction;}
  void GenericAddrSpace(bool genericAddrSpace) {_genericAddrSpace = genericAddrSpace;}
  bool GenericAddrSpace() {return _genericAddrSpace;}
  void OldValueCheck(bool check) {_oldValueCheck = check;}
  bool OldValueCheck() {return _oldValueCheck;}
  void LocalRefValues(bool localRefValues) {_localRefValues = localRefValues;}
  bool LocalRefValues() {return _localRefValues;}
  void MaxGroupSize(cl_uint maxGroupSize) {_maxGroupSize = maxGroupSize;}
  cl_uint MaxGroupSize() {return _maxGroupSize;}
  void CurrentGroupSize(cl_uint currentGroupSize)
  {
    if(MaxGroupSize() && MaxGroupSize() < currentGroupSize)
      _currentGroupSize = MaxGroupSize();
    else
      _currentGroupSize = currentGroupSize;
  }
  cl_uint CurrentGroupSize() {return _currentGroupSize;}
  virtual cl_uint CurrentGroupNum(cl_uint threadCount)
  {
    if(threadCount == 0)
      return 0;
    if(LocalMemory())
      return 1;
    return threadCount/CurrentGroupSize();
  }
  cl_int Iterations() {return _iterations;}
  std::string IterationsStr() {std::stringstream ss; ss << _iterations; return ss.str();}
private:
  const TExplicitAtomicType _dataType;
  const bool _useSVM;
  HostDataType	_startValue;
  bool _localMemory;
  bool _declaredInProgram;
  bool _usedInFunction;
  bool _genericAddrSpace;
  bool _oldValueCheck;
  bool _localRefValues;
  cl_uint _maxGroupSize;
  cl_uint _currentGroupSize;
  cl_uint _passCount;
  const cl_int _iterations;
};

template<typename HostAtomicType, typename HostDataType>
class CBasicTestMemOrderScope : public CBasicTest<HostAtomicType, HostDataType>
{
public:
  using CBasicTest<HostAtomicType, HostDataType>::LocalMemory;
  using CBasicTest<HostAtomicType, HostDataType>::MaxGroupSize;
  using CBasicTest<HostAtomicType, HostDataType>::CheckCapabilities;
  CBasicTestMemOrderScope(TExplicitAtomicType dataType, bool useSVM = false) : CBasicTest<HostAtomicType, HostDataType>(dataType, useSVM)
  {
  }
  virtual std::string ProgramHeader(cl_uint maxNumDestItems)
  {
    std::string header;
    if(gOldAPI)
    {
      std::string s = MemoryScope() == MEMORY_SCOPE_EMPTY ? "" : ",s";
      header +=
        "#define atomic_store_explicit(x,y,o"+s+")                     atomic_store(x,y)\n"
        "#define atomic_load_explicit(x,o"+s+")                        atomic_load(x)\n"
        "#define atomic_exchange_explicit(x,y,o"+s+")                  atomic_exchange(x,y)\n"
        "#define atomic_compare_exchange_strong_explicit(x,y,z,os,of"+s+") atomic_compare_exchange_strong(x,y,z)\n"
        "#define atomic_compare_exchange_weak_explicit(x,y,z,os,of"+s+")   atomic_compare_exchange_weak(x,y,z)\n"
        "#define atomic_fetch_add_explicit(x,y,o"+s+")                 atomic_fetch_add(x,y)\n"
        "#define atomic_fetch_sub_explicit(x,y,o"+s+")                 atomic_fetch_sub(x,y)\n"
        "#define atomic_fetch_or_explicit(x,y,o"+s+")                  atomic_fetch_or(x,y)\n"
        "#define atomic_fetch_xor_explicit(x,y,o"+s+")                 atomic_fetch_xor(x,y)\n"
        "#define atomic_fetch_and_explicit(x,y,o"+s+")                 atomic_fetch_and(x,y)\n"
        "#define atomic_fetch_min_explicit(x,y,o"+s+")                 atomic_fetch_min(x,y)\n"
        "#define atomic_fetch_max_explicit(x,y,o"+s+")                 atomic_fetch_max(x,y)\n"
        "#define atomic_flag_test_and_set_explicit(x,o"+s+")           atomic_flag_test_and_set(x)\n"
        "#define atomic_flag_clear_explicit(x,o"+s+")                  atomic_flag_clear(x)\n";
    }
    return header+CBasicTest<HostAtomicType, HostDataType>::ProgramHeader(maxNumDestItems);
  }
  virtual std::string SingleTestName()
  {
    std::string testName = CBasicTest<HostAtomicType, HostDataType>::SingleTestName();
    if(MemoryOrder() != MEMORY_ORDER_EMPTY)
    {
      testName += std::string(", ")+std::string(get_memory_order_type_name(MemoryOrder())).substr(sizeof("memory"));
    }
    if(MemoryScope() != MEMORY_SCOPE_EMPTY)
    {
      testName += std::string(", ")+std::string(get_memory_scope_type_name(MemoryScope())).substr(sizeof("memory"));
    }
    return testName;
  }
  virtual int ExecuteSingleTest(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    if(LocalMemory() &&
      MemoryScope() != MEMORY_SCOPE_EMPTY &&
      MemoryScope() != MEMORY_SCOPE_WORK_GROUP) //memory scope should only be used for global memory
      return 0;
    if(MemoryScope() == MEMORY_SCOPE_DEVICE)
      MaxGroupSize(16); // increase number of groups by forcing smaller group size
    else
      MaxGroupSize(0); // group size limited by device capabilities

    if (CheckCapabilities(MemoryScope(), MemoryOrder()) == TEST_SKIPPED_ITSELF)
        return 0; // skip test - not applicable

    return CBasicTest<HostAtomicType, HostDataType>::ExecuteSingleTest(deviceID, context, queue);
  }
  virtual int ExecuteForEachParameterSet(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    // repeat test for each reasonable memory order/scope combination
    std::vector<TExplicitMemoryOrderType> memoryOrder;
    std::vector<TExplicitMemoryScopeType> memoryScope;
    int error = 0;

    // For OpenCL-3.0 and later some orderings and scopes are optional, so here
    // we query for the supported ones.
    test_error_ret(
        getSupportedMemoryOrdersAndScopes(deviceID, memoryOrder, memoryScope),
        "getSupportedMemoryOrdersAndScopes failed\n", TEST_FAIL);

    for(unsigned oi = 0; oi < memoryOrder.size(); oi++)
    {
      for(unsigned si = 0; si < memoryScope.size(); si++)
      {
        if(memoryOrder[oi] == MEMORY_ORDER_EMPTY && memoryScope[si] != MEMORY_SCOPE_EMPTY)
          continue;
        MemoryOrder(memoryOrder[oi]);
        MemoryScope(memoryScope[si]);
        EXECUTE_TEST(error, (CBasicTest<HostAtomicType, HostDataType>::ExecuteForEachParameterSet(deviceID, context, queue)));
      }
    }
    return error;
  }
  void MemoryOrder(TExplicitMemoryOrderType memoryOrder) {_memoryOrder = memoryOrder;}
  TExplicitMemoryOrderType MemoryOrder() {return _memoryOrder;}
  std::string MemoryOrderStr()
  {
    if(MemoryOrder() != MEMORY_ORDER_EMPTY)
      return std::string(", ")+get_memory_order_type_name(MemoryOrder());
    return "";
  }
  void MemoryScope(TExplicitMemoryScopeType memoryScope) {_memoryScope = memoryScope;}
  TExplicitMemoryScopeType MemoryScope() {return _memoryScope;}
  std::string MemoryScopeStr()
  {
    if(MemoryScope() != MEMORY_SCOPE_EMPTY)
      return std::string(", ")+get_memory_scope_type_name(MemoryScope());
    return "";
  }
  std::string MemoryOrderScopeStr()
  {
    return MemoryOrderStr()+MemoryScopeStr();
  }
  virtual cl_uint CurrentGroupNum(cl_uint threadCount)
  {
    if(MemoryScope() == MEMORY_SCOPE_WORK_GROUP)
      return 1;
    return CBasicTest<HostAtomicType, HostDataType>::CurrentGroupNum(threadCount);
  }
  virtual cl_uint MaxHostThreads()
  {
      // block host threads execution for memory scope different than
      // memory_scope_all_svm_devices
      if (MemoryScope() == MEMORY_SCOPE_ALL_DEVICES
          || MemoryScope() == MEMORY_SCOPE_ALL_SVM_DEVICES || gHost)
      {
          return CBasicTest<HostAtomicType, HostDataType>::MaxHostThreads();
      }
      else
      {
          return 0;
      }
  }
private:
  TExplicitMemoryOrderType _memoryOrder;
  TExplicitMemoryScopeType _memoryScope;
};

template<typename HostAtomicType, typename HostDataType>
class CBasicTestMemOrder2Scope : public CBasicTestMemOrderScope<HostAtomicType, HostDataType>
{
public:
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::LocalMemory;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrder;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScope;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryOrderStr;
  using CBasicTestMemOrderScope<HostAtomicType, HostDataType>::MemoryScopeStr;
  using CBasicTest<HostAtomicType, HostDataType>::CheckCapabilities;

  CBasicTestMemOrder2Scope(TExplicitAtomicType dataType, bool useSVM = false) : CBasicTestMemOrderScope<HostAtomicType, HostDataType>(dataType, useSVM)
  {
  }
  virtual std::string SingleTestName()
  {
    std::string testName = CBasicTest<HostAtomicType, HostDataType>::SingleTestName();
    if(MemoryOrder() != MEMORY_ORDER_EMPTY)
      testName += std::string(", ")+std::string(get_memory_order_type_name(MemoryOrder())).substr(sizeof("memory"));
    if(MemoryOrder2() != MEMORY_ORDER_EMPTY)
      testName += std::string(", ")+std::string(get_memory_order_type_name(MemoryOrder2())).substr(sizeof("memory"));
    if(MemoryScope() != MEMORY_SCOPE_EMPTY)
      testName += std::string(", ")+std::string(get_memory_scope_type_name(MemoryScope())).substr(sizeof("memory"));
    return testName;
  }
  virtual int ExecuteForEachParameterSet(cl_device_id deviceID, cl_context context, cl_command_queue queue)
  {
    // repeat test for each reasonable memory order/scope combination
    std::vector<TExplicitMemoryOrderType> memoryOrder;
    std::vector<TExplicitMemoryScopeType> memoryScope;
    int error = 0;

    // For OpenCL-3.0 and later some orderings and scopes are optional, so here
    // we query for the supported ones.
    test_error_ret(
        getSupportedMemoryOrdersAndScopes(deviceID, memoryOrder, memoryScope),
        "getSupportedMemoryOrdersAndScopes failed\n", TEST_FAIL);

    for(unsigned oi = 0; oi < memoryOrder.size(); oi++)
    {
      for(unsigned o2i = 0; o2i < memoryOrder.size(); o2i++)
      {
        for(unsigned si = 0; si < memoryScope.size(); si++)
        {
          if((memoryOrder[oi] == MEMORY_ORDER_EMPTY || memoryOrder[o2i] == MEMORY_ORDER_EMPTY)
            && memoryOrder[oi] != memoryOrder[o2i])
            continue; // both memory order arguments must be set (or none)
          if((memoryOrder[oi] == MEMORY_ORDER_EMPTY || memoryOrder[o2i] == MEMORY_ORDER_EMPTY)
            && memoryScope[si] != MEMORY_SCOPE_EMPTY)
            continue; // memory scope without memory order is not allowed
          MemoryOrder(memoryOrder[oi]);
          MemoryOrder2(memoryOrder[o2i]);
          MemoryScope(memoryScope[si]);

          if (CheckCapabilities(MemoryScope(), MemoryOrder())
              == TEST_SKIPPED_ITSELF)
              continue; // skip test - not applicable

          if (CheckCapabilities(MemoryScope(), MemoryOrder2())
              == TEST_SKIPPED_ITSELF)
              continue; // skip test - not applicable

          EXECUTE_TEST(error, (CBasicTest<HostAtomicType, HostDataType>::ExecuteForEachParameterSet(deviceID, context, queue)));
        }
      }
    }
    return error;
  }
  void MemoryOrder2(TExplicitMemoryOrderType memoryOrderFail) {_memoryOrder2 = memoryOrderFail;}
  TExplicitMemoryOrderType MemoryOrder2() {return _memoryOrder2;}
  std::string MemoryOrderFailStr()
  {
    if(MemoryOrder2() != MEMORY_ORDER_EMPTY)
      return std::string(", ")+get_memory_order_type_name(MemoryOrder2());
    return "";
  }
  std::string MemoryOrderScope()
  {
    return MemoryOrderStr()+MemoryOrderFailStr()+MemoryScopeStr();
  }
private:
  TExplicitMemoryOrderType _memoryOrder2;
};

template<typename HostAtomicType, typename HostDataType>
std::string CBasicTest<HostAtomicType, HostDataType>::PragmaHeader(cl_device_id deviceID)
{
  std::string pragma;

  if(gOldAPI)
  {
    pragma += "#pragma OPENCL EXTENSION cl_khr_local_int32_base_atomics : enable\n";
    pragma += "#pragma OPENCL EXTENSION cl_khr_local_int32_extended_atomics : enable\n";
    pragma += "#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : enable\n";
    pragma += "#pragma OPENCL EXTENSION cl_khr_global_int32_extended_atomics : enable\n";
  }
  // Create the pragma lines for this kernel
  if(DataType().Size(deviceID) == 8)
  {
    pragma += "#pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable\n";
    pragma += "#pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable\n";
  }
  if(_dataType == TYPE_ATOMIC_DOUBLE)
    pragma += "#pragma OPENCL EXTENSION cl_khr_fp64 : enable\n";
  return pragma;
}

template<typename HostAtomicType, typename HostDataType>
std::string CBasicTest<HostAtomicType, HostDataType>::ProgramHeader(cl_uint maxNumDestItems)
{
  // Create the program header
  std::string header;
  std::string aTypeName = DataType().AtomicTypeName();
  std::string cTypeName = DataType().RegularTypeName();
  std::string argListForKernel;
  std::string argListForFunction;
  std::string argListNoTypes;
  std::string functionPrototype;
  std::string addressSpace = LocalMemory() ? "__local " : "__global ";

  if(gOldAPI)
  {
    header += std::string("#define ")+aTypeName+" "+cTypeName+"\n"
      "#define atomic_store(x,y)                                (*(x) = y)\n"
      "#define atomic_load(x)                                   (*(x))\n"
      "#define ATOMIC_VAR_INIT(x)                               (x)\n"
      "#define ATOMIC_FLAG_INIT                                 0\n"
      "#define atomic_init(x,y)                                 atomic_store(x,y)\n";
    if(aTypeName == "atomic_float")
      header += "#define atomic_exchange(x,y)                             atomic_xchg(x,y)\n";
    else if(aTypeName == "atomic_double")
      header += "double atomic_exchange(volatile "+addressSpace+"atomic_double *x, double y)\n"
        "{\n"
        "  long tmp = *(long*)&y, res;\n"
        "  volatile "+addressSpace+"long *tmpA = (volatile "+addressSpace+"long)x;\n"
        "  res = atom_xchg(tmpA,tmp);\n"
        "  return *(double*)&res;\n"
        "}\n";
    else
      header += "#define atomic_exchange(x,y)                             atom_xchg(x,y)\n";
    if(aTypeName != "atomic_float" && aTypeName != "atomic_double")
      header +=
      "bool atomic_compare_exchange_strong(volatile "+addressSpace+" "+aTypeName+" *a, "+cTypeName+" *expected, "+cTypeName+" desired)\n"
      "{\n"
      "  "+cTypeName+" old = atom_cmpxchg(a, *expected, desired);\n"
      "  if(old == *expected)\n"
      "    return true;\n"
      "  *expected = old;\n"
      "  return false;\n"
      "}\n"
      "#define atomic_compare_exchange_weak                     atomic_compare_exchange_strong\n";
    header +=
      "#define atomic_fetch_add(x,y)                            atom_add(x,y)\n"
      "#define atomic_fetch_sub(x,y)                            atom_sub(x,y)\n"
      "#define atomic_fetch_or(x,y)                             atom_or(x,y)\n"
      "#define atomic_fetch_xor(x,y)                            atom_xor(x,y)\n"
      "#define atomic_fetch_and(x,y)                            atom_and(x,y)\n"
      "#define atomic_fetch_min(x,y)                            atom_min(x,y)\n"
      "#define atomic_fetch_max(x,y)                            atom_max(x,y)\n"
      "#define atomic_flag_test_and_set(x)                      atomic_exchange(x,1)\n"
      "#define atomic_flag_clear(x)                             atomic_store(x,0)\n"
      "\n";
  }
  if(!LocalMemory() && DeclaredInProgram())
  {
    // additional atomic variable for results copying (last thread will do this)
    header += "__global volatile atomic_uint finishedThreads = ATOMIC_VAR_INIT(0);\n";
    // atomic variables declared in program scope - test data
    std::stringstream ss;
    ss << maxNumDestItems;
    header += std::string("__global volatile ")+aTypeName+" destMemory["+ss.str()+"] = {\n";
    ss.str("");
    ss << _startValue;
    for(cl_uint i = 0; i < maxNumDestItems; i++)
    {
      if(aTypeName == "atomic_flag")
        header +=  "  ATOMIC_FLAG_INIT";
      else
        header +=  "  ATOMIC_VAR_INIT("+ss.str()+")";
      if(i+1 < maxNumDestItems)
        header += ",";
      header += "\n";
    }
    header+=
      "};\n"
      "\n";
  }
  return header;
}

template<typename HostAtomicType, typename HostDataType>
std::string CBasicTest<HostAtomicType, HostDataType>::FunctionCode()
{
  if(!UsedInFunction())
    return "";
  std::string addressSpace = LocalMemory() ? "__local " : "__global ";
  std::string code = "void test_atomic_function(uint tid, uint threadCount, uint numDestItems, volatile ";
  if(!GenericAddrSpace())
    code += addressSpace;
  code += std::string(DataType().AtomicTypeName())+" *destMemory, __global "+DataType().RegularTypeName()+
    " *oldValues";
  if(LocalRefValues())
    code += std::string(", __local ")+DataType().RegularTypeName()+" *localValues";
  code += ")\n"
    "{\n";
  code += ProgramCore();
  code += "}\n"
    "\n";
  return code;
}

template<typename HostAtomicType, typename HostDataType>
std::string CBasicTest<HostAtomicType, HostDataType>::KernelCode(cl_uint maxNumDestItems)
{
  std::string aTypeName = DataType().AtomicTypeName();
  std::string cTypeName = DataType().RegularTypeName();
  std::string addressSpace = LocalMemory() ? "__local " : "__global ";
  std::string code = "__kernel void test_atomic_kernel(uint threadCount, uint numDestItems, ";

  // prepare list of arguments for kernel
  if(LocalMemory())
  {
    code += std::string("__global ")+cTypeName+" *finalDest, __global "+cTypeName+" *oldValues,"
      " volatile "+addressSpace+aTypeName+" *"+(DeclaredInProgram() ? "notUsed" : "")+"destMemory";
  }
  else
  {
    code += "volatile "+addressSpace+(DeclaredInProgram() ? (cTypeName+" *finalDest") : (aTypeName+" *destMemory"))+
      ", __global "+cTypeName+" *oldValues";
  }
  if(LocalRefValues())
    code += std::string(", __local ")+cTypeName+" *localValues";
  code += ")\n"
    "{\n";
  if(LocalMemory() && DeclaredInProgram())
  {
    // local atomics declared in kernel scope
    std::stringstream ss;
    ss << maxNumDestItems;
    code += std::string("  __local volatile ")+aTypeName+" destMemory["+ss.str()+"];\n";
  }
  code += "  uint  tid = get_global_id(0);\n"
    "\n";
  if(LocalMemory())
  {
      // memory_order_relaxed is sufficient for these initialization operations
      // as the barrier below will act as a fence, providing an order to the
      // operations. memory_scope_work_group is sufficient as local memory is
      // only visible within the work-group.
      code += R"(
              // initialize atomics not reachable from host (first thread
              // is doing this, other threads are waiting on barrier)
              if(get_local_id(0) == 0)
                for(uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)
                {)";
      if (aTypeName == "atomic_flag")
      {
          code += R"(
                  if(finalDest[dstItemIdx])
                    atomic_flag_test_and_set_explicit(destMemory+dstItemIdx,
                                                      memory_order_relaxed,
                                                      memory_scope_work_group);
                  else
                    atomic_flag_clear_explicit(destMemory+dstItemIdx,
                                               memory_order_relaxed,
                                               memory_scope_work_group);)";
      }
    else
    {
        code += R"(
                atomic_store_explicit(destMemory+dstItemIdx,
                                      finalDest[dstItemIdx],
                                      memory_order_relaxed,
                                      memory_scope_work_group);)";
    }
    code +=
      "    }\n"
      "  barrier(CLK_LOCAL_MEM_FENCE);\n"
      "\n";
  }
  if (LocalRefValues())
  {
    code +=
      "  // Copy input reference values into local memory\n";
    if (NumNonAtomicVariablesPerThread() == 1)
      code += "  localValues[get_local_id(0)] = oldValues[tid];\n";
    else
    {
      std::stringstream ss;
      ss << NumNonAtomicVariablesPerThread();
      code +=
        "  for(uint rfId = 0; rfId < " + ss.str() + "; rfId++)\n"
        "    localValues[get_local_id(0)*" + ss.str() + "+rfId] = oldValues[tid*" + ss.str() + "+rfId];\n";
    }
    code +=
      "  barrier(CLK_LOCAL_MEM_FENCE);\n"
      "\n";
  }
  if (UsedInFunction())
    code += std::string("  test_atomic_function(tid, threadCount, numDestItems, destMemory, oldValues")+
    (LocalRefValues() ? ", localValues" : "")+");\n";
  else
    code += ProgramCore();
  code += "\n";
  if (LocalRefValues())
  {
    code +=
      "  // Copy local reference values into output array\n"
      "  barrier(CLK_LOCAL_MEM_FENCE);\n";
    if (NumNonAtomicVariablesPerThread() == 1)
      code += "  oldValues[tid] = localValues[get_local_id(0)];\n";
    else
    {
      std::stringstream ss;
      ss << NumNonAtomicVariablesPerThread();
      code +=
        "  for(uint rfId = 0; rfId < " + ss.str() + "; rfId++)\n"
        "    oldValues[tid*" + ss.str() + "+rfId] = localValues[get_local_id(0)*" + ss.str() + "+rfId];\n";
    }
    code += "\n";
  }
  if(LocalMemory() || DeclaredInProgram())
  {
    code += "  // Copy final values to host reachable buffer\n";
    if(LocalMemory())
      code +=
        "  barrier(CLK_LOCAL_MEM_FENCE);\n"
        "  if(get_local_id(0) == 0) // first thread in workgroup\n";
    else
      // global atomics declared in program scope
      code += R"(
                if(atomic_fetch_add_explicit(&finishedThreads, 1u,
                                           memory_order_relaxed,
                                           memory_scope_work_group)
                   == get_global_size(0)-1) // last finished thread
                   )";
    code +=
        "    for(uint dstItemIdx = 0; dstItemIdx < numDestItems; dstItemIdx++)\n";
    if(aTypeName == "atomic_flag")
    {
        code += R"(
                finalDest[dstItemIdx] =
                    atomic_flag_test_and_set_explicit(destMemory+dstItemIdx,
                                                      memory_order_relaxed,
                                                      memory_scope_work_group);)";
    }
    else
    {
        code += R"(
                finalDest[dstItemIdx] =
                    atomic_load_explicit(destMemory+dstItemIdx,
                                         memory_order_relaxed,
                                         memory_scope_work_group);)";
    }
  }
  code += "}\n"
    "\n";
  return code;
}

template <typename HostAtomicType, typename HostDataType>
int CBasicTest<HostAtomicType, HostDataType>::ExecuteSingleTest(cl_device_id deviceID, cl_context context, cl_command_queue queue)
{
  int error;
  clProgramWrapper program;
  clKernelWrapper kernel;
  size_t threadNum[1];
  clMemWrapper streams[2];
  std::vector<HostAtomicType> destItems;
  HostAtomicType *svmAtomicBuffer = 0;
  std::vector<HostDataType> refValues, startRefValues;
  HostDataType *svmDataBuffer = 0;
  cl_uint deviceThreadCount, hostThreadCount, threadCount;
  size_t groupSize = 0;
  std::string programSource;
  const char *programLine;
  MTdata d;
  size_t typeSize = DataType().Size(deviceID);

  deviceThreadCount = _maxDeviceThreads;
  hostThreadCount = MaxHostThreads();
  threadCount = deviceThreadCount+hostThreadCount;

  //log_info("\t%s %s%s...\n", local ? "local" : "global", DataType().AtomicTypeName(), memoryOrderScope.c_str());
  log_info("\t%s...\n", SingleTestName().c_str());

  if(!LocalMemory() && DeclaredInProgram() && gNoGlobalVariables) // no support for program scope global variables
  {
    log_info("\t\tTest disabled\n");
    return 0;
  }
  if(UsedInFunction() && GenericAddrSpace() && gNoGenericAddressSpace)
  {
    log_info("\t\tTest disabled\n");
    return 0;
  }

  // set up work sizes based on device capabilities and test configuration
  error = clGetDeviceInfo(deviceID, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(groupSize), &groupSize, NULL);
  test_error(error, "Unable to obtain max work group size for device");
  CurrentGroupSize((cl_uint)groupSize);
  if(CurrentGroupSize() > deviceThreadCount)
    CurrentGroupSize(deviceThreadCount);
  if(CurrentGroupNum(deviceThreadCount) == 1 || gOldAPI)
    deviceThreadCount = CurrentGroupSize()*CurrentGroupNum(deviceThreadCount);
  threadCount = deviceThreadCount+hostThreadCount;

  // If we're given a num_results function, we need to determine how many result objects we need.
  // This is the first assessment for current maximum number of threads (exact thread count is not known here)
  // - needed for program source code generation (arrays of atomics declared in program)
  cl_uint numDestItems = NumResults(threadCount, deviceID);

  if(deviceThreadCount > 0)
  {
      // This loop iteratively reduces the workgroup size by 2 and then
      // re-generates the kernel with the reduced
      // workgroup size until we find a size which is admissible for the kernel
      // being run or reduce the wg size
      // to the trivial case of 1 (which was separately verified to be accurate
      // for the kernel being run)

      while ((CurrentGroupSize() > 1))
      {
          // Re-generate the kernel code with the current group size
          if (kernel) clReleaseKernel(kernel);
          if (program) clReleaseProgram(program);
          programSource = PragmaHeader(deviceID) + ProgramHeader(numDestItems)
              + FunctionCode() + KernelCode(numDestItems);
          programLine = programSource.c_str();
          if (create_single_kernel_helper_with_build_options(
                  context, &program, &kernel, 1, &programLine,
                  "test_atomic_kernel", gOldAPI ? "" : nullptr))
          {
              return -1;
          }
          // Get work group size for the new kernel
          error = clGetKernelWorkGroupInfo(kernel, deviceID,
                                           CL_KERNEL_WORK_GROUP_SIZE,
                                           sizeof(groupSize), &groupSize, NULL);
          test_error(error,
                     "Unable to obtain max work group size for device and "
                     "kernel combo");

          if (LocalMemory())
          {
              cl_ulong usedLocalMemory;
              cl_ulong totalLocalMemory;
              cl_uint maxWorkGroupSize;

              error = clGetKernelWorkGroupInfo(
                  kernel, deviceID, CL_KERNEL_LOCAL_MEM_SIZE,
                  sizeof(usedLocalMemory), &usedLocalMemory, NULL);
              test_error(error, "clGetKernelWorkGroupInfo failed");

              error = clGetDeviceInfo(deviceID, CL_DEVICE_LOCAL_MEM_SIZE,
                                      sizeof(totalLocalMemory),
                                      &totalLocalMemory, NULL);
              test_error(error, "clGetDeviceInfo failed");

              // We know that each work-group is going to use typeSize *
              // deviceThreadCount bytes of local memory
              // so pick the maximum value for deviceThreadCount that uses all
              // the local memory.
              maxWorkGroupSize =
                  ((totalLocalMemory - usedLocalMemory) / typeSize);

              if (maxWorkGroupSize < groupSize) groupSize = maxWorkGroupSize;
          }
          if (CurrentGroupSize() <= groupSize)
              break;
          else
              CurrentGroupSize(CurrentGroupSize() / 2);
      }
    if(CurrentGroupSize() > deviceThreadCount)
      CurrentGroupSize(deviceThreadCount);
    if(CurrentGroupNum(deviceThreadCount) == 1 || gOldAPI)
      deviceThreadCount = CurrentGroupSize()*CurrentGroupNum(deviceThreadCount);
    threadCount = deviceThreadCount+hostThreadCount;
  }
  if (gDebug)
  {
      log_info("Program source:\n");
      log_info("%s\n", programLine);
  }
  if(deviceThreadCount > 0)
    log_info("\t\t(thread count %u, group size %u)\n", deviceThreadCount, CurrentGroupSize());
  if(hostThreadCount > 0)
    log_info("\t\t(host threads %u)\n", hostThreadCount);

  refValues.resize(threadCount*NumNonAtomicVariablesPerThread());

  // Generate ref data if we have a ref generator provided
  d = init_genrand(gRandomSeed);
  startRefValues.resize(threadCount*NumNonAtomicVariablesPerThread());
  if(GenerateRefs(threadCount, &startRefValues[0], d))
  {
    //copy ref values for host threads
    memcpy(&refValues[0], &startRefValues[0], sizeof(HostDataType)*threadCount*NumNonAtomicVariablesPerThread());
  }
  else
  {
    startRefValues.resize(0);
  }
  free_mtdata(d);
  d = NULL;

  // If we're given a num_results function, we need to determine how many result objects we need. If
  // we don't have it, we assume it's just 1
  // This is final value (exact thread count is known in this place)
  numDestItems = NumResults(threadCount, deviceID);

  destItems.resize(numDestItems);
  for(cl_uint i = 0; i < numDestItems; i++)
    destItems[i] = _startValue;

  // Create main buffer with atomic variables (array size dependent on particular test)
  if(UseSVM())
  {
    if(gUseHostPtr)
      svmAtomicBuffer = (HostAtomicType*)malloc(typeSize * numDestItems);
    else
      svmAtomicBuffer = (HostAtomicType*)clSVMAlloc(context, CL_MEM_SVM_FINE_GRAIN_BUFFER | CL_MEM_SVM_ATOMICS, typeSize * numDestItems, 0);
    if(!svmAtomicBuffer)
    {
      log_error("ERROR: clSVMAlloc failed!\n");
      return -1;
    }
    memcpy(svmAtomicBuffer, &destItems[0], typeSize * numDestItems);
    streams[0] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
                                typeSize * numDestItems, svmAtomicBuffer, NULL);
  }
  else
  {
      streams[0] = clCreateBuffer(context, CL_MEM_COPY_HOST_PTR,
                                  typeSize * numDestItems, &destItems[0], NULL);
  }
  if (!streams[0])
  {
    log_error("ERROR: Creating output array failed!\n");
    return -1;
  }
  // Create buffer for per-thread input/output data
  if(UseSVM())
  {
    if(gUseHostPtr)
      svmDataBuffer = (HostDataType*)malloc(typeSize*threadCount*NumNonAtomicVariablesPerThread());
    else
      svmDataBuffer = (HostDataType*)clSVMAlloc(context, CL_MEM_SVM_FINE_GRAIN_BUFFER | (SVMDataBufferAllSVMConsistent() ? CL_MEM_SVM_ATOMICS : 0), typeSize*threadCount*NumNonAtomicVariablesPerThread(), 0);
    if(!svmDataBuffer)
    {
      log_error("ERROR: clSVMAlloc failed!\n");
      return -1;
    }
    if(startRefValues.size())
      memcpy(svmDataBuffer, &startRefValues[0], typeSize*threadCount*NumNonAtomicVariablesPerThread());
    streams[1] = clCreateBuffer(context, CL_MEM_USE_HOST_PTR,
                                typeSize * threadCount
                                    * NumNonAtomicVariablesPerThread(),
                                svmDataBuffer, NULL);
  }
  else
  {
      streams[1] = clCreateBuffer(
          context,
          ((startRefValues.size() ? CL_MEM_COPY_HOST_PTR : CL_MEM_READ_WRITE)),
          typeSize * threadCount * NumNonAtomicVariablesPerThread(),
          startRefValues.size() ? &startRefValues[0] : 0, NULL);
  }
  if (!streams[1])
  {
    log_error("ERROR: Creating reference array failed!\n");
    return -1;
  }
  if(deviceThreadCount > 0)
  {
    cl_uint argInd = 0;
    /* Set the arguments */
    error = clSetKernelArg(kernel, argInd++, sizeof(threadCount), &threadCount);
    test_error(error, "Unable to set kernel argument");
    error = clSetKernelArg(kernel, argInd++, sizeof(numDestItems), &numDestItems);
    test_error(error, "Unable to set indexed kernel argument");
    error = clSetKernelArg(kernel, argInd++, sizeof(streams[0]), &streams[0]);
    test_error(error, "Unable to set indexed kernel arguments");
    error = clSetKernelArg(kernel, argInd++, sizeof(streams[1]), &streams[1]);
    test_error(error, "Unable to set indexed kernel arguments");
    if(LocalMemory())
    {
      error = clSetKernelArg(kernel, argInd++, typeSize * numDestItems, NULL);
      test_error(error, "Unable to set indexed local kernel argument");
    }
    if(LocalRefValues())
    {
      error = clSetKernelArg(kernel, argInd++, LocalRefValues() ? typeSize*CurrentGroupSize()*NumNonAtomicVariablesPerThread() : 1, NULL);
      test_error(error, "Unable to set indexed kernel argument");
    }
  }
  /* Configure host threads */
  std::vector<THostThreadContext> hostThreadContexts(hostThreadCount);
  for(unsigned int t = 0; t < hostThreadCount; t++)
  {
    hostThreadContexts[t].test = this;
    hostThreadContexts[t].tid = deviceThreadCount+t;
    hostThreadContexts[t].threadCount = threadCount;
    hostThreadContexts[t].destMemory = UseSVM() ? svmAtomicBuffer : &destItems[0];
    hostThreadContexts[t].oldValues = UseSVM() ? svmDataBuffer : &refValues[0];
  }

  if(deviceThreadCount > 0)
  {
    /* Run the kernel */
    threadNum[0] = deviceThreadCount;
    groupSize = CurrentGroupSize();
    error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threadNum, &groupSize, 0, NULL, NULL);
    test_error(error, "Unable to execute test kernel");
    /* start device threads */
    error = clFlush(queue);
    test_error(error, "clFlush failed");
  }

  /* Start host threads and wait for finish */
  if(hostThreadCount > 0)
    ThreadPool_Do(HostThreadFunction, hostThreadCount, &hostThreadContexts[0]);

  if(UseSVM())
  {
    error = clFinish(queue);
    test_error(error, "clFinish failed");
    memcpy(&destItems[0], svmAtomicBuffer, typeSize*numDestItems);
    memcpy(&refValues[0], svmDataBuffer, typeSize*threadCount*NumNonAtomicVariablesPerThread());
  }
  else
  {
    if(deviceThreadCount > 0)
    {
      error = clEnqueueReadBuffer(queue, streams[0], CL_TRUE, 0, typeSize * numDestItems, &destItems[0], 0, NULL, NULL);
      test_error(error, "Unable to read result value!");
      error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, typeSize * deviceThreadCount*NumNonAtomicVariablesPerThread(), &refValues[0], 0, NULL, NULL);
      test_error(error, "Unable to read reference values!");
    }
  }
  bool dataVerified = false;
  // If we have an expectedFn, then we need to generate a final value to compare against. If we don't
  // have one, it's because we're comparing ref values only
  for(cl_uint i = 0; i < numDestItems; i++)
  {
    HostDataType expected;

    if(!ExpectedValue(expected, threadCount, startRefValues.size() ? &startRefValues[0] : 0, i))
      break; // no expected value function provided

    if(expected != destItems[i])
    {
      std::stringstream logLine;
      logLine << "ERROR: Result " << i << " from kernel does not validate! (should be " << expected << ", was " << destItems[i] << ")\n";
      log_error("%s", logLine.str().c_str());
      for(i = 0; i < threadCount; i++)
      {
        logLine.str("");
        logLine << " --- " << i << " - ";
        if(startRefValues.size())
          logLine << startRefValues[i] << " -> " << refValues[i];
        else
          logLine << refValues[i];
        logLine << " --- ";
        if(i < numDestItems)
          logLine << destItems[i];
        logLine << "\n";
        log_info("%s", logLine.str().c_str());
      }
      if(!gDebug)
      {
        log_info("Program source:\n");
        log_info("%s\n", programLine);
      }
      return -1;
    }
    dataVerified = true;
  }

  bool dataCorrect = false;
  /* Use the verify function (if provided) to also check the results */
  if(VerifyRefs(dataCorrect, threadCount, &refValues[0], &destItems[0]))
  {
    if(!dataCorrect)
    {
      log_error("ERROR: Reference values did not validate!\n");
      std::stringstream logLine;
      for(cl_uint i = 0; i < threadCount; i++)
      for (cl_uint j = 0; j < NumNonAtomicVariablesPerThread(); j++)
      {
        logLine.str("");
        logLine << " --- " << i << " - " << refValues[i*NumNonAtomicVariablesPerThread()+j] << " --- ";
        if(j == 0 && i < numDestItems)
          logLine << destItems[i];
        logLine << "\n";
        log_info("%s", logLine.str().c_str());
      }
      if(!gDebug)
      {
        log_info("Program source:\n");
        log_info("%s\n", programLine);
      }
      return -1;
    }
  }
  else if(!dataVerified)
  {
    log_error("ERROR: Test doesn't check total or refs; no values are verified!\n");
    return -1;
  }

  if(OldValueCheck() &&
    !(DeclaredInProgram() && !LocalMemory())) // don't test for programs scope global atomics
                                             // 'old' value has been overwritten by previous clEnqueueNDRangeKernel
  {
    /* Re-write the starting value */
    for(size_t i = 0; i < numDestItems; i++)
      destItems[i] = _startValue;
    refValues[0] = 0;
    if(deviceThreadCount > 0)
    {
      error = clEnqueueWriteBuffer(queue, streams[0], CL_TRUE, 0, typeSize * numDestItems, &destItems[0], 0, NULL, NULL);
      test_error(error, "Unable to write starting values!");

      /* Run the kernel once for a single thread, so we can verify that the returned value is the original one */
      threadNum[0] = 1;
      error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, threadNum, threadNum, 0, NULL, NULL);
      test_error(error, "Unable to execute test kernel");

      error = clEnqueueReadBuffer(queue, streams[1], CL_TRUE, 0, typeSize, &refValues[0], 0, NULL, NULL);
      test_error(error, "Unable to read reference values!");
    }
    else
    {
      /* Start host thread */
      HostFunction(0, 1, &destItems[0], &refValues[0]);
    }

    if(refValues[0] != _startValue)//destItems[0])
    {
      std::stringstream logLine;
      logLine << "ERROR: atomic function operated correctly but did NOT return correct 'old' value "
        " (should have been " << destItems[0] << ", returned " << refValues[0] << ")!\n";
      log_error("%s", logLine.str().c_str());
      if(!gDebug)
      {
        log_info("Program source:\n");
        log_info("%s\n", programLine);
      }
      return -1;
    }
  }
  if(UseSVM())
  {
    // the buffer object must first be released before the SVM buffer is freed
    error = clReleaseMemObject(streams[0]);
    streams[0] = 0;
    test_error(error, "clReleaseMemObject failed");
    if(gUseHostPtr)
      free(svmAtomicBuffer);
    else
      clSVMFree(context, svmAtomicBuffer);
    error = clReleaseMemObject(streams[1]);
    streams[1] = 0;
    test_error(error, "clReleaseMemObject failed");
    if(gUseHostPtr)
      free(svmDataBuffer);
    else
      clSVMFree(context, svmDataBuffer);
  }
  _passCount++;
  return 0;
}

#endif //_COMMON_H_