Skip to content

Commit 816b8e5

Browse files
ktfktf
authored
Merge dfd9e31 into sapling-pr-archive-ktf
2 parents 9a5f0e0 + dfd9e31 commit 816b8e5

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

47 files changed

+484
-387
lines changed

‎Common/MathUtils/include/MathUtils/SMatrixGPU.h‎

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -29,6 +29,7 @@
2929
#include "GPUCommonMath.h"
3030
#include "GPUCommonAlgorithm.h"
3131
#include "GPUCommonLogger.h"
32+
#include "GPUCommonTypeTraits.h"
3233

3334
namespace o2::math_utils::detail
3435
{
@@ -468,6 +469,9 @@ class SMatrixGPU
468469
GPUd() const T& operator()(unsigned int i, unsigned int j) const;
469470
GPUd() T& operator()(unsigned int i, unsigned int j);
470471

472+
template <typename Y, typename X>
473+
GPUd() friend X& operator<<(Y& y, const SMatrixGPU&);
474+
471475
class SMatrixRowGPU
472476
{
473477
public:
@@ -512,6 +516,15 @@ class SMatrixGPU
512516
R mRep;
513517
};
514518

519+
#ifndef __OPENCL__ // TODO: current C++ for OpenCL 2021 is at C++17, so no concepts. But we don't need this trick for OpenCL anyway, so we can just hide it.
520+
template <class T, unsigned int D1, unsigned int D2, class R, typename Y, typename X = Y>
521+
requires(sizeof(typename X::traits_type::pos_type) != 0) // do not provide a template to fair::Logger, etc... (pos_type is a member type of all std::ostream classes)
522+
GPUd() X& operator<<(Y& y, const SMatrixGPU<T, D1, D2, R>&)
523+
{
524+
return y;
525+
}
526+
#endif
527+
515528
template <class T, unsigned int D1, unsigned int D2, class R>
516529
GPUdi() SMatrixGPU<T, D1, D2, R>::SMatrixGPU(SMatrixIdentity)
517530
{

‎DataFormats/Detectors/TRD/src/Tracklet64.cxx‎

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -40,6 +40,16 @@ std::ostream& operator<<(std::ostream& stream, const Tracklet64& trg)
4040
trg.printStream(stream);
4141
return stream;
4242
}
43+
44+
bool operator<(const Tracklet64& lhs, const Tracklet64& rhs)
45+
{
46+
return (lhs.getDetector() < rhs.getDetector()) ||
47+
(lhs.getDetector() == rhs.getDetector() && lhs.getROB() < rhs.getROB()) ||
48+
(lhs.getDetector() == rhs.getDetector() && lhs.getROB() == rhs.getROB() && lhs.getMCM() < rhs.getMCM()) ||
49+
(lhs.getDetector() == rhs.getDetector() && lhs.getROB() == rhs.getROB() && lhs.getMCM() == rhs.getMCM() && lhs.getPadRow() < rhs.getPadRow()) ||
50+
(lhs.getDetector() == rhs.getDetector() && lhs.getROB() == rhs.getROB() && lhs.getMCM() == rhs.getMCM() && lhs.getPadRow() == rhs.getPadRow() && lhs.getPadCol() < rhs.getPadCol());
51+
}
52+
4353
#endif // GPUCA_GPUCODE_DEVICE
4454

4555
} // namespace trd

‎DataFormats/simulation/include/SimulationDataFormat/O2DatabasePDG.h‎

Lines changed: 4 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -500,6 +500,10 @@ inline void O2DatabasePDG::addALICEParticles(TDatabasePDG* db)
500500
}
501501

502502
// glueball hunting
503+
ionCode = 115;
504+
if (!db->GetParticle(ionCode)) {
505+
db->AddParticle("a2_1320", "a2_1320", 1.3182, kFALSE, 0.1078, 0, "Resonance", ionCode);
506+
}
503507
ionCode = 10221;
504508
if (!db->GetParticle(ionCode)) {
505509
db->AddParticle("f0_1370", "f0_1370", 1.37, kFALSE, 0.200, 0, "Resonance", ionCode);

‎Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -842,14 +842,14 @@ void countCellsHandler(
842842
cellsLUTsHost, // d_in
843843
cellsLUTsHost, // d_out
844844
nTracklets + 1, // num_items
845-
0));
845+
0)); // NOLINT: this is the offset of the sum, not a pointer
846846
discardResult(cudaMalloc(&d_temp_storage, temp_storage_bytes));
847847
gpuCheckError(cub::DeviceScan::ExclusiveSum(d_temp_storage, // d_temp_storage
848848
temp_storage_bytes, // temp_storage_bytes
849849
cellsLUTsHost, // d_in
850850
cellsLUTsHost, // d_out
851851
nTracklets + 1, // num_items
852-
0));
852+
0)); // NOLINT: this is the offset of the sum, not a pointer
853853
// gpu::printBufferLayerOnThread<<<1, 1>>>(layer, cellsLUTsHost, nTracklets + 1);
854854
gpuCheckError(cudaFree(d_temp_storage));
855855
}
@@ -934,14 +934,14 @@ void countCellNeighboursHandler(CellSeed** cellsLayersDevice,
934934
neighboursIndexTable, // d_in
935935
neighboursIndexTable, // d_out
936936
nCells + 1, // num_items
937-
0));
937+
0)); // NOLINT: this is the offset of the sum, not a pointer
938938
discardResult(cudaMalloc(&d_temp_storage_2, temp_storage_bytes_2));
939939
gpuCheckError(cub::DeviceScan::ExclusiveSum(d_temp_storage_2, // d_temp_storage
940940
temp_storage_bytes_2, // temp_storage_bytes
941941
neighboursIndexTable, // d_in
942942
neighboursIndexTable, // d_out
943943
nCells + 1, // num_items
944-
0));
944+
0)); // NOLINT: this is the offset of the sum, not a pointer
945945
gpuCheckError(cudaFree(d_temp_storage));
946946
gpuCheckError(cudaFree(d_temp_storage_2));
947947
gpuCheckError(cudaPeekAtLastError());
@@ -1032,4 +1032,4 @@ void trackSeedHandler(CellSeed* trackSeeds,
10321032
gpuCheckError(cudaPeekAtLastError());
10331033
gpuCheckError(cudaDeviceSynchronize());
10341034
}
1035-
} // namespace o2::its
1035+
} // namespace o2::its

‎Framework/Core/include/Framework/TableTreeHelpers.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -105,7 +105,7 @@ class ColumnToBranch
105105
arrow::Type::type mFieldType;
106106
std::vector<uint8_t> cache;
107107
std::shared_ptr<arrow::Array> mCurrentArray = nullptr;
108-
int64_t mChunkLength;
108+
int64_t mChunkLength = 0;
109109
int mFieldSize = 0;
110110
};
111111

‎GPU/Common/GPUCommonConstants.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -17,7 +17,7 @@
1717

1818
#include "GPUCommonDef.h"
1919

20-
#if !defined(__OPENCL__) || defined(__OPENCLCPP__)
20+
#if !defined(__OPENCL1__)
2121
namespace GPUCA_NAMESPACE::gpu::gpu_common_constants
2222
{
2323
static CONSTEXPR const float kCLight = 0.000299792458f;

‎GPU/Common/GPUCommonDef.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,7 @@
3030
//Some GPU configuration settings, must be included first
3131
#include "GPUCommonDefSettings.h"
3232

33-
#if (!defined(__OPENCL__) || defined(__OPENCLCPP__)) && (!(defined(__CINT__) || defined(__ROOTCINT__)) || defined(__CLING__)) && defined(__cplusplus) && __cplusplus >= 201103L
33+
#if !defined(__OPENCL1__) && (!(defined(__CINT__) || defined(__ROOTCINT__)) || defined(__CLING__)) && defined(__cplusplus) && __cplusplus >= 201103L
3434
#define GPUCA_NOCOMPAT // C++11 + No old ROOT5 + No old OpenCL
3535
#ifndef __OPENCL__
3636
#define GPUCA_NOCOMPAT_ALLOPENCL // + No OpenCL at all
@@ -82,7 +82,7 @@
8282
#define GPUCA_NAMESPACE o2
8383
#endif
8484

85-
#if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCL__) && !defined(__OPENCLCPP__) && defined(GPUCA_OPENCL_NO_CONSTANT_MEMORY)) || (defined(__OPENCLCPP__) && defined(GPUCA_OPENCLCPP_NO_CONSTANT_MEMORY))
85+
#if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCL1__) && defined(GPUCA_OPENCL_NO_CONSTANT_MEMORY)) || (defined(__OPENCLCPP__) && defined(GPUCA_OPENCLCPP_NO_CONSTANT_MEMORY))
8686
#define GPUCA_NO_CONSTANT_MEMORY
8787
#elif defined(__CUDACC__) || defined(__HIPCC__)
8888
#define GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM

‎GPU/Common/GPUCommonMath.h‎

Lines changed: 5 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -31,7 +31,7 @@
3131
#include <cstdint>
3232
#endif
3333

34-
#if !defined(__OPENCL__) || defined(__OPENCLCPP__)
34+
#if !defined(__OPENCL1__)
3535
namespace GPUCA_NAMESPACE
3636
{
3737
namespace gpu
@@ -220,7 +220,7 @@ GPUdi() uint32_t GPUCommonMath::Float2UIntReint(const float& x)
220220
{
221221
#if defined(GPUCA_GPUCODE_DEVICE) && (defined(__CUDACC__) || defined(__HIPCC__))
222222
return __float_as_uint(x);
223-
#elif defined(GPUCA_GPUCODE_DEVICE) && (defined(__OPENCL__) || defined(__OPENCLCPP__))
223+
#elif defined(GPUCA_GPUCODE_DEVICE) && defined(__OPENCL__)
224224
return as_uint(x);
225225
#else
226226
return reinterpret_cast<const uint32_t&>(x);
@@ -289,7 +289,7 @@ GPUhdi() void GPUCommonMath::SinCosd(double x, double& s, double& c)
289289

290290
GPUdi() uint32_t GPUCommonMath::Clz(uint32_t x)
291291
{
292-
#if (defined(__GNUC__) || defined(__clang__) || defined(__CUDACC__) || defined(__HIPCC__)) && (!defined(__OPENCL__) || defined(__OPENCLCPP__))
292+
#if (defined(__GNUC__) || defined(__clang__) || defined(__CUDACC__) || defined(__HIPCC__)) && !defined(__OPENCL1__)
293293
return x == 0 ? 32 : CHOICE(__builtin_clz(x), __clz(x), __builtin_clz(x)); // use builtin if available
294294
#else
295295
for (int32_t i = 31; i >= 0; i--) {
@@ -303,7 +303,7 @@ GPUdi() uint32_t GPUCommonMath::Clz(uint32_t x)
303303

304304
GPUdi() uint32_t GPUCommonMath::Popcount(uint32_t x)
305305
{
306-
#if (defined(__GNUC__) || defined(__clang__) || defined(__CUDACC__) || defined(__HIPCC__)) && (!defined(__OPENCL__) /*|| defined(__OPENCLCPP__)*/) // TODO: remove OPENCLCPP workaround when reported SPIR-V bug is fixed
306+
#if (defined(__GNUC__) || defined(__clang__) || defined(__CUDACC__) || defined(__HIPCC__)) && (!defined(__OPENCL__) /* !defined(__OPENCL1__)*/) // TODO: exclude only OPENCLC (not CPP) when reported SPIR-V bug is fixed
307307
// use builtin if available
308308
return CHOICE(__builtin_popcount(x), __popc(x), __builtin_popcount(x));
309309
#else
@@ -563,7 +563,7 @@ GPUdii() void GPUCommonMath::AtomicMinInternal(GPUglobalref() GPUgeneric() GPUAt
563563

564564
#undef CHOICE
565565

566-
#if !defined(__OPENCL__) || defined(__OPENCLCPP__)
566+
#if !defined(__OPENCL1__)
567567
}
568568
}
569569
#endif

‎GPU/Common/GPUCommonRtypes.h‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -20,14 +20,14 @@
2020
#if defined(GPUCA_STANDALONE) || (defined(GPUCA_O2_LIB) && !defined(GPUCA_O2_INTERFACE)) || defined(GPUCA_GPUCODE) // clang-format off
2121
#if !defined(ROOT_Rtypes) && !defined(__CLING__)
2222
#define GPUCOMMONRTYPES_H_ACTIVE
23+
struct MUST_NOT_USE_Rtypes_h {};
24+
typedef MUST_NOT_USE_Rtypes_h TClass;
2325
#define ClassDef(name,id)
2426
#define ClassDefNV(name, id)
2527
#define ClassDefOverride(name, id)
2628
#define ClassImp(name)
2729
#define templateClassImp(name)
2830
#ifndef GPUCA_GPUCODE_DEVICE
29-
// typedef uint64_t ULong64_t;
30-
// typedef uint32_t UInt_t;
3131
#include <iostream>
3232
#endif
3333
#endif

‎GPU/Common/GPUCommonTypeTraits.h‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -21,7 +21,7 @@
2121
#ifndef GPUCA_GPUCODE_COMPILEKERNELS
2222
#include <type_traits>
2323
#endif
24-
#elif !defined(__OPENCL__) || defined(__OPENCLCPP__)
24+
#elif !defined(__OPENCL1__)
2525
// We just reimplement some type traits in std for the GPU
2626
namespace std
2727
{

0 commit comments

Comments
 (0)