Project
Loading...
Searching...
No Matches
GPUDef.h
Go to the documentation of this file.
1// Copyright 2019-2020 CERN and copyright holders of ALICE O2.
2// See https://alice-o2.web.cern.ch/copyright for details of the copyright holders.
3// All rights not expressly granted are reserved.
4//
5// This software is distributed under the terms of the GNU General Public
6// License v3 (GPL Version 3), copied verbatim in the file "COPYING".
7//
8// In applying this license CERN does not waive the privileges and immunities
9// granted to it by virtue of its status as an Intergovernmental Organization
10// or submit itself to any jurisdiction.
11
14
15// clang-format off
16#ifndef GPUDEF_H
17#define GPUDEF_H
18
19#include "GPUCommonDef.h"
22#include "GPUCommonRtypes.h"
23
24// Macros for masking ptrs in OpenCL kernel calls as uint64_t (The API only allows us to pass buffer objects)
25#ifdef __OPENCL__
26 #define GPUPtr1(a, b) uint64_t b
27 #ifdef __OPENCL__
28 #define GPUPtr2(a, b) ((__generic a) (a) b)
29 #else
30 #define GPUPtr2(a, b) ((__global a) (a) b)
31 #endif
32#else
33 #define GPUPtr1(a, b) a b
34 #define GPUPtr2(a, b) b
35#endif
36
37#define GPUCA_EVDUMP_FILE "event"
38
39#ifdef GPUCA_GPUCODE
40 #define CA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared) const GPUsharedref() vartype& __restrict__ varname = varshared;
41 #define CA_SHARED_STORAGE(storage) storage
42 #define CA_SHARED_CACHE(target, src, size) \
43 static_assert((size) % sizeof(int32_t) == 0, "Invalid shared cache size"); \
44 for (uint32_t i_shared_cache = get_local_id(0); i_shared_cache < (size) / sizeof(int32_t); i_shared_cache += get_local_size(0)) { \
45 reinterpret_cast<GPUsharedref() int32_t*>(target)[i_shared_cache] = reinterpret_cast<GPUglobalref() const int32_t*>(src)[i_shared_cache]; \
46 }
47 #define CA_SHARED_CACHE_REF(target, src, size, reftype, ref) \
48 CA_SHARED_CACHE(target, src, size) \
49 GPUsharedref() const reftype* __restrict__ ref = (target)
50#else
51 #define CA_MAKE_SHARED_REF(vartype, varname, varglobal, varshared) const GPUglobalref() vartype & __restrict__ varname = varglobal;
52 #define CA_SHARED_STORAGE(storage)
53 #define CA_SHARED_CACHE(target, src, size)
54 #define CA_SHARED_CACHE_REF(target, src, size, reftype, ref) GPUglobalref() const reftype* __restrict__ ref = src
55#endif
56
57#endif //GPUTPCDEF_H
58
59#ifdef GPUCA_CADEBUG
60 #ifdef CADEBUG
61 #undef CADEBUG
62 #endif
63 #ifdef GPUCA_CADEBUG_ENABLED
64 #undef GPUCA_CADEBUG_ENABLED
65 #endif
66 #if GPUCA_CADEBUG == 1 && !defined(GPUCA_GPUCODE)
67 #define CADEBUG(...) __VA_ARGS__
68 #define CADEBUG2(cmd, ...) {__VA_ARGS__; cmd;}
69 #define GPUCA_CADEBUG_ENABLED
70 #endif
71 #undef GPUCA_CADEBUG
72#endif
73
74#ifndef CADEBUG
75 #define CADEBUG(...)
76 #define CADEBUG2(cmd, ...) {cmd;}
77#endif
78// clang-format on