Skip to content

Please consider the following formatting changes to #13865 #109

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
alibuild wants to merge 4 commits into from
Closed

Please consider the following formatting changes to #13865 #109

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
e01c675
GPU:Remove OpenCL 1.2, now that OpenCL 2 supports all its functionality
davidrohr Jan 15, 2025
514e953
GPU: Remove some template magic that was only needed for OpenCL 1.2
davidrohr Jan 15, 2025
9e155cf
Code-Checker, silence some false warnings from failures in clang-tidy
davidrohr Jan 15, 2025
e2888c4
Please consider the following formatting changes
alibuild Jan 15, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
4 changes: 2 additions & 2 deletions Detectors/ITSMFT/ITS/tracking/GPU/cuda/TrackingKernels.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1215,14 +1215,14 @@ void processNeighboursHandler(const int startLayer,
thrust::raw_pointer_cast(&foundSeedsTable[0]), // d_in
thrust::raw_pointer_cast(&foundSeedsTable[0]), // d_out
nCurrentCells + 1, // num_items
0));
0)); // NOLINT: failure in clang-tidy
discardResult(cudaMalloc(&d_temp_storage, temp_storage_bytes));
gpuCheckError(cub::DeviceScan::ExclusiveSum(d_temp_storage, // d_temp_storage
temp_storage_bytes, // temp_storage_bytes
thrust::raw_pointer_cast(&foundSeedsTable[0]), // d_in
thrust::raw_pointer_cast(&foundSeedsTable[0]), // d_out
nCurrentCells + 1, // num_items
0));
0)); // NOLINT: failure in clang-tidy

thrust::device_vector<int> updatedCellIds(foundSeedsTable.back()) /*, lastCellIds(foundSeedsTable.back())*/;
thrust::device_vector<CellSeed> updatedCellSeeds(foundSeedsTable.back()) /*, lastCellSeeds(foundSeedsTable.back())*/;
Expand Down
4 changes: 1 addition & 3 deletions GPU/Common/GPUCommonConstants.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,11 +17,9 @@

#include "GPUCommonDef.h"

#if !defined(__OPENCL1__)
namespace GPUCA_NAMESPACE::gpu::gpu_common_constants
{
static CONSTEXPR const float kCLight = 0.000299792458f; // TODO: Duplicate of MathConstants, fix this when OpenCL1 is removed
static CONSTEXPR const float kCLight = 0.000299792458f; // TODO: Duplicate of MathConstants, fix this now that we use only OpenCL CPP
}
#endif

#endif
4 changes: 2 additions & 2 deletions GPU/Common/GPUCommonDef.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,7 @@
//Some GPU configuration settings, must be included first
#include "GPUCommonDefSettings.h"

#if !defined(__OPENCL1__) && (!(defined(__CINT__) || defined(__ROOTCINT__)) || defined(__CLING__)) && defined(__cplusplus) && __cplusplus >= 201103L
#if (!(defined(__CINT__) || defined(__ROOTCINT__)) || defined(__CLING__)) && defined(__cplusplus) && __cplusplus >= 201103L
#define GPUCA_NOCOMPAT // C++11 + No old ROOT5 + No old OpenCL
#ifndef __OPENCL__
#define GPUCA_NOCOMPAT_ALLOPENCL // + No OpenCL at all
Expand Down Expand Up @@ -82,7 +82,7 @@
#define GPUCA_NAMESPACE o2
#endif

#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))
#if (defined(__CUDACC__) && defined(GPUCA_CUDA_NO_CONSTANT_MEMORY)) || (defined(__HIPCC__) && defined(GPUCA_HIP_NO_CONSTANT_MEMORY)) || (defined(__OPENCLCPP__) && defined(GPUCA_OPENCLCPP_NO_CONSTANT_MEMORY))
#define GPUCA_NO_CONSTANT_MEMORY
#elif defined(__CUDACC__) || defined(__HIPCC__)
#define GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM
Expand Down
1 change: 0 additions & 1 deletion GPU/Common/GPUCommonDefSettings.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -26,7 +26,6 @@

//#define GPUCA_CUDA_NO_CONSTANT_MEMORY // Do not use constant memory for CUDA
//#define GPUCA_HIP_NO_CONSTANT_MEMORY // Do not use constant memory for HIP
//#define GPUCA_OPENCL_NO_CONSTANT_MEMORY // Do not use constant memory for OpenCL 1.2
#define GPUCA_OPENCLCPP_NO_CONSTANT_MEMORY // Do not use constant memory for OpenCL C++ - MANDATORY as OpenCL cannot cast between __constant and __generic yet!

// clang-format on
Expand Down
12 changes: 4 additions & 8 deletions GPU/Common/GPUCommonMath.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -31,12 +31,10 @@
#include <cstdint>
#endif

#if !defined(__OPENCL1__)
namespace GPUCA_NAMESPACE
{
namespace gpu
{
#endif

class GPUCommonMath
{
Expand Down Expand Up @@ -289,7 +287,7 @@ GPUhdi() void GPUCommonMath::SinCosd(double x, double& s, double& c)

GPUdi() uint32_t GPUCommonMath::Clz(uint32_t x)
{
#if (defined(__GNUC__) || defined(__clang__) || defined(__CUDACC__) || defined(__HIPCC__)) && !defined(__OPENCL1__)
#if (defined(__GNUC__) || defined(__clang__) || defined(__CUDACC__) || defined(__HIPCC__))
return x == 0 ? 32 : CHOICE(__builtin_clz(x), __clz(x), __builtin_clz(x)); // use builtin if available
#else
for (int32_t i = 31; i >= 0; i--) {
Expand All @@ -303,7 +301,7 @@ GPUdi() uint32_t GPUCommonMath::Clz(uint32_t x)

GPUdi() uint32_t GPUCommonMath::Popcount(uint32_t x)
{
#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
#if (defined(__GNUC__) || defined(__clang__) || defined(__CUDACC__) || defined(__HIPCC__)) && !defined(__OPENCL__) // TODO: remove OPENCL when reported SPIR-V bug is fixed
// use builtin if available
return CHOICE(__builtin_popcount(x), __popc(x), __builtin_popcount(x));
#else
Expand Down Expand Up @@ -563,9 +561,7 @@ GPUdii() void GPUCommonMath::AtomicMinInternal(GPUglobalref() GPUgeneric() GPUAt

#undef CHOICE

#if !defined(__OPENCL1__)
}
}
#endif
} // namespace gpu
} // namespace GPUCA_NAMESPACE

#endif // GPUCOMMONMATH_H
2 changes: 1 addition & 1 deletion GPU/Common/GPUCommonTypeTraits.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,7 +21,7 @@
#ifndef GPUCA_GPUCODE_COMPILEKERNELS
#include <type_traits>
#endif
#elif !defined(__OPENCL1__)
#else
// We just reimplement some type traits in std for the GPU
namespace std
{
Expand Down
14 changes: 6 additions & 8 deletions GPU/GPUTracking/Base/GPUConstantMem.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,7 +20,7 @@
#include "GPUDataTypes.h"
#include "GPUErrors.h"

// Dummies for stuff not supported in legacy code (ROOT 5 / OPENCL1.2)
// Dummies for stuff not supported in legacy code (ROOT 5)
#if defined(GPUCA_NOCOMPAT_ALLCINT)
#include "GPUTPCGMMerger.h"
#else
Expand Down Expand Up @@ -71,12 +71,10 @@ namespace GPUCA_NAMESPACE
{
namespace gpu
{
MEM_CLASS_PRE()
struct GPUConstantMem {
MEM_CONSTANT(GPUParam)
param;
MEM_GLOBAL(GPUTPCTracker)
tpcTrackers[GPUCA_NSLICES];
GPUParam param;
GPUTPCTracker
tpcTrackers[GPUCA_NSLICES];
GPUTPCConvert tpcConverter;
GPUTPCCompression tpcCompressor;
GPUTPCDecompression tpcDecompressor;
Expand Down Expand Up @@ -150,7 +148,7 @@ namespace gpu
{

// Must be placed here, to avoid circular header dependency
GPUdi() GPUconstantref() const MEM_CONSTANT(GPUConstantMem) * GPUProcessor::GetConstantMem() const
GPUdi() GPUconstantref() const GPUConstantMem* GPUProcessor::GetConstantMem() const
{
#if defined(GPUCA_GPUCODE_DEVICE) && defined(GPUCA_HAS_GLOBAL_SYMBOL_CONSTANT_MEM) && !defined(GPUCA_GPUCODE_HOSTONLY)
return &GPUCA_CONSMEM;
Expand All @@ -159,7 +157,7 @@ GPUdi() GPUconstantref() const MEM_CONSTANT(GPUConstantMem) * GPUProcessor::GetC
#endif
}

GPUdi() GPUconstantref() const MEM_CONSTANT(GPUParam) & GPUProcessor::Param() const
GPUdi() GPUconstantref() const GPUParam& GPUProcessor::Param() const
{
return GetConstantMem()->param;
}
Expand Down
4 changes: 2 additions & 2 deletions GPU/GPUTracking/Base/GPUGeneralKernels.cxx
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,7 +17,7 @@
using namespace GPUCA_NAMESPACE::gpu;

template <>
GPUdii() void GPUMemClean16::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& GPUrestrict() processors, GPUglobalref() void* ptr, uint64_t size)
GPUdii() void GPUMemClean16::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors, GPUglobalref() void* ptr, uint64_t size)
{
const uint64_t stride = get_global_size(0);
int4 i0;
Expand All @@ -30,7 +30,7 @@ GPUdii() void GPUMemClean16::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_
}

template <>
GPUdii() void GPUitoa::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& GPUrestrict() processors, GPUglobalref() int32_t* ptr, uint64_t size)
GPUdii() void GPUitoa::Thread<0>(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& GPUrestrict() processors, GPUglobalref() int32_t* ptr, uint64_t size)
{
const uint64_t stride = get_global_size(0);
for (uint64_t i = get_global_id(0); i < size; i += stride) {
Expand Down
15 changes: 6 additions & 9 deletions GPU/GPUTracking/Base/GPUGeneralKernels.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,6 @@ namespace GPUCA_NAMESPACE
{
namespace gpu
{
MEM_CLASS_PRE()
struct GPUConstantMem;

class GPUKernelTemplate
Expand All @@ -50,7 +49,6 @@ class GPUKernelTemplate
step4 = 4,
step5 = 5 };

MEM_CLASS_PRE()
struct GPUSharedMemory {
};

Expand Down Expand Up @@ -82,21 +80,20 @@ class GPUKernelTemplate
#endif
};

typedef GPUconstantref() MEM_CONSTANT(GPUConstantMem) processorType;
typedef GPUconstantref() GPUConstantMem processorType;
GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep() { return GPUCA_RECO_STEP::NoRecoStep; }
MEM_TEMPLATE()
GPUhdi() static processorType* Processor(MEM_TYPE(GPUConstantMem) & processors)
GPUhdi() static processorType* Processor(GPUConstantMem& processors)
{
return &processors;
}
#ifdef GPUCA_NOCOMPAT
template <int32_t iKernel, typename... Args>
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& processors, Args... args)
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, Args... args)
{
}
#else
template <int32_t iKernel>
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& processors)
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors)
{
}
#endif
Expand All @@ -108,7 +105,7 @@ class GPUMemClean16 : public GPUKernelTemplate
public:
GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep() { return GPUCA_RECO_STEP::NoRecoStep; }
template <int32_t iKernel = defaultKernel>
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& processors, GPUglobalref() void* ptr, uint64_t size);
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUglobalref() void* ptr, uint64_t size);
};

// Fill with incrementing sequnce of integers
Expand All @@ -117,7 +114,7 @@ class GPUitoa : public GPUKernelTemplate
public:
GPUhdi() CONSTEXPR static GPUDataTypes::RecoStep GetRecoStep() { return GPUCA_RECO_STEP::NoRecoStep; }
template <int32_t iKernel = defaultKernel>
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() MEM_LOCAL(GPUSharedMemory) & smem, processorType& processors, GPUglobalref() int32_t* ptr, uint64_t size);
GPUd() static void Thread(int32_t nBlocks, int32_t nThreads, int32_t iBlock, int32_t iThread, GPUsharedref() GPUSharedMemory& smem, processorType& processors, GPUglobalref() int32_t* ptr, uint64_t size);
};

} // namespace gpu
Expand Down
1 change: 0 additions & 1 deletion GPU/GPUTracking/Base/GPUParam.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -79,7 +79,6 @@ struct GPUParam_t {
} // namespace internal

#if !(defined(__CINT__) || defined(__ROOTCINT__)) || defined(__CLING__) // Hide from ROOT 5 CINT
MEM_CLASS_PRE()
struct GPUParam : public internal::GPUParam_t<GPUSettingsRec, GPUSettingsParam> {

#ifndef GPUCA_GPUCODE
Expand Down
Loading
Loading