GPUDef.h

// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
// All rights not expressly granted are reserved.
//
// This software is distributed under the terms of the GNU General Public
// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
//
// In applying this license CERN does not waive the privileges and immunities
// granted to it by virtue of its status as an Intergovernmental Organization
// or submit itself to any jurisdiction.

/// \file GPUDef.h
/// \author David Rohr, Sergey Gorbunov

// clang-format off
#ifndef GPUDEF_H
#define GPUDEF_H

#include "GPUCommonDef.h"
#include "GPUDefConstantsAndSettings.h"
#include "GPUDefParametersWrapper.h"
#include "GPUCommonRtypes.h"

// Macros for masking ptrs in OpenCL kernel calls as uint64_t (The API only allows us to pass buffer objects)
#ifdef __OPENCL__
  #define GPUPtr1(a, b) uint64_t b
  #ifdef __OPENCL__
    #define GPUPtr2(a, b) ((__generic a) (a) b)
  #else
    #define GPUPtr2(a, b) ((__global a) (a) b)
  #endif
#else
  #define GPUPtr1(a, b) a b
  #define GPUPtr2(a, b) b
#endif

#define GPUCA_EVDUMP_FILE "event"

#ifdef GPUCA_GPUCODE
  #define CA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared) const GPUsharedref() vartype& __restrict__ varname = varshared;
  #define CA_SHARED_STORAGE(storage) storage
  #define CA_SHARED_CACHE(target, src, size) \
    static_assert((size) % sizeof(int32_t) == 0, "Invalid shared cache size"); \
    for (uint32_t i_shared_cache = get_local_id(0); i_shared_cache < (size) / sizeof(int32_t); i_shared_cache += get_local_size(0)) { \
      reinterpret_cast<GPUsharedref() int32_t*>(target)[i_shared_cache] = reinterpret_cast<GPUglobalref() const int32_t*>(src)[i_shared_cache]; \
    }
  #define CA_SHARED_CACHE_REF(target, src, size, reftype, ref) \
    CA_SHARED_CACHE(target, src, size) \
    GPUsharedref() const reftype* __restrict__ ref = (target)
#else
  #define CA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared) const GPUglobalref() vartype & __restrict__ varname = varglobal;
  #define CA_SHARED_STORAGE(storage)
  #define CA_SHARED_CACHE(target, src, size)
  #define CA_SHARED_CACHE_REF(target, src, size, reftype, ref) GPUglobalref() const reftype* __restrict__ ref = src
#endif

#endif //GPUTPCDEF_H

#ifdef GPUCA_CADEBUG
  #ifdef CADEBUG
    #undef CADEBUG
  #endif
  #if GPUCA_CADEBUG == 1 && !defined(GPUCA_GPUCODE)
    #define CADEBUG(...) __VA_ARGS__
    #define CADEBUG2(cmd, ...) {__VA_ARGS__; cmd;}
    #define GPUCA_CADEBUG_ENABLED
  #endif
  #undef GPUCA_CADEBUG
#endif

#ifndef CADEBUG
  #define CADEBUG(...)
  #define CADEBUG2(cmd, ...) {cmd;}
#endif
// clang-format on