diff --git a/Detectors/ITSMFT/ITS/tracking/cuda/include/ITStrackingCUDA/PrimaryVertexContextNV.h b/Detectors/ITSMFT/ITS/tracking/cuda/include/ITStrackingCUDA/PrimaryVertexContextNV.h index 5eb5b2a459d33..9809b36bebbcb 100644 --- a/Detectors/ITSMFT/ITS/tracking/cuda/include/ITStrackingCUDA/PrimaryVertexContextNV.h +++ b/Detectors/ITSMFT/ITS/tracking/cuda/include/ITStrackingCUDA/PrimaryVertexContextNV.h @@ -36,7 +36,7 @@ class PrimaryVertexContextNV final : public PrimaryVertexContext { public: PrimaryVertexContextNV() = default; - virtual ~PrimaryVertexContextNV() = default; + ~PrimaryVertexContextNV() override; void initialise(const MemoryParameters& memParam, const TrackingParameters& trkParam, const std::vector>& cl, const std::array& pv, const int iteration) override; @@ -62,6 +62,8 @@ class PrimaryVertexContextNV final : public PrimaryVertexContext std::array, constants::its2::CellsPerRoad - 1> mTempCellArray; }; +inline PrimaryVertexContextNV::~PrimaryVertexContextNV() = default; + inline gpu::DeviceStoreNV& PrimaryVertexContextNV::getDeviceContext() { return *mGPUContextDevicePointer; diff --git a/Detectors/ITSMFT/ITS/tracking/cuda/include/ITStrackingCUDA/VertexerTraitsGPU.h b/Detectors/ITSMFT/ITS/tracking/cuda/include/ITStrackingCUDA/VertexerTraitsGPU.h index 8ae1b578cc261..48c3c7043a598 100644 --- a/Detectors/ITSMFT/ITS/tracking/cuda/include/ITStrackingCUDA/VertexerTraitsGPU.h +++ b/Detectors/ITSMFT/ITS/tracking/cuda/include/ITStrackingCUDA/VertexerTraitsGPU.h @@ -47,7 +47,7 @@ class VertexerTraitsGPU : public VertexerTraits ~VertexerTraitsGPU() override; #else VertexerTraitsGPU(); - ~VertexerTraitsGPU() = default; + ~VertexerTraitsGPU() override; #endif void initialise(ROframe*) override; void computeTracklets() override; diff --git a/Detectors/ITSMFT/ITS/tracking/cuda/src/VertexerTraitsGPU.cu b/Detectors/ITSMFT/ITS/tracking/cuda/src/VertexerTraitsGPU.cu index 336465f6a8ddf..db3d58b2b14bf 100644 --- a/Detectors/ITSMFT/ITS/tracking/cuda/src/VertexerTraitsGPU.cu +++ b/Detectors/ITSMFT/ITS/tracking/cuda/src/VertexerTraitsGPU.cu @@ -93,6 +93,8 @@ VertexerTraitsGPU::VertexerTraitsGPU() #endif +VertexerTraitsGPU::~VertexerTraitsGPU() = default; + void VertexerTraitsGPU::initialise(ROframe* event) { reset(); diff --git a/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/UniquePointerHIP.h b/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/UniquePointerHIP.h index 372e2ee31f19f..afe6e22af7a8b 100644 --- a/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/UniquePointerHIP.h +++ b/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/UniquePointerHIP.h @@ -82,8 +82,8 @@ UniquePointer::UniquePointer(const T& ref) { try { - Utils::HostHIP::gpuMalloc(reinterpret_cast(&mDevicePointer), sizeof(T)); - Utils::HostHIP::gpuMemcpyHostToDevice(mDevicePointer, &ref, sizeof(T)); + utils::host_hip::gpuMalloc(reinterpret_cast(&mDevicePointer), sizeof(T)); + utils::host_hip::gpuMemcpyHostToDevice(mDevicePointer, &ref, sizeof(T)); } catch (...) { @@ -119,7 +119,7 @@ void UniquePointer::destroy() { if (mDevicePointer != nullptr) { - Utils::HostHIP::gpuFree(mDevicePointer); + utils::host_hip::gpuFree(mDevicePointer); } } diff --git a/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/UtilsHIP.h b/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/UtilsHIP.h index 34434e9ec242c..f1d80e8e7b56b 100644 --- a/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/UtilsHIP.h +++ b/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/UtilsHIP.h @@ -26,10 +26,10 @@ namespace its namespace gpu { -namespace Utils +namespace utils { -namespace HostHIP +namespace host_hip { #ifdef __HIPCC__ @@ -50,15 +50,15 @@ void gpuMemcpyHostToDeviceAsync(void*, const void*, int, hipStream_t&); void gpuMemcpyDeviceToHost(void*, const void*, int); // void gpuStartProfiler(); // void gpuStopProfiler(); -} // namespace Host +} // namespace host_hip // -namespace DeviceHIP +namespace device_hip { GPUd() int getLaneIndex(); GPUd() int shareToWarp(const int, const int); GPUd() int gpuAtomicAdd(int*, const int); -} // namespace Device -} // namespace Utils +} // namespace device_hip +} // namespace utils } // namespace gpu } // namespace its } // namespace o2 diff --git a/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/VectorHIP.h b/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/VectorHIP.h index d7aaa7c3b2a4a..9f15ac3ed430c 100644 --- a/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/VectorHIP.h +++ b/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/VectorHIP.h @@ -100,17 +100,17 @@ VectorHIP::VectorHIP(const T* const source, const int size, const int initial if (size > 0) { try { - Utils::HostHIP::gpuMalloc(reinterpret_cast(&mArrayPointer), size * sizeof(T)); - Utils::HostHIP::gpuMalloc(reinterpret_cast(&mDeviceSize), sizeof(int)); + utils::host_hip::gpuMalloc(reinterpret_cast(&mArrayPointer), size * sizeof(T)); + utils::host_hip::gpuMalloc(reinterpret_cast(&mDeviceSize), sizeof(int)); if (source != nullptr) { - Utils::HostHIP::gpuMemcpyHostToDevice(mArrayPointer, source, size * sizeof(T)); - Utils::HostHIP::gpuMemcpyHostToDevice(mDeviceSize, &size, sizeof(int)); + utils::host_hip::gpuMemcpyHostToDevice(mArrayPointer, source, size * sizeof(T)); + utils::host_hip::gpuMemcpyHostToDevice(mDeviceSize, &size, sizeof(int)); } else { - Utils::HostHIP::gpuMemcpyHostToDevice(mDeviceSize, &initialSize, sizeof(int)); + utils::host_hip::gpuMemcpyHostToDevice(mDeviceSize, &initialSize, sizeof(int)); } } catch (...) { @@ -179,7 +179,7 @@ template int VectorHIP::getSizeFromDevice() const { int size; - Utils::HostHIP::gpuMemcpyDeviceToHost(&size, mDeviceSize, sizeof(int)); + utils::host_hip::gpuMemcpyDeviceToHost(&size, mDeviceSize, sizeof(int)); return size; } @@ -187,7 +187,7 @@ int VectorHIP::getSizeFromDevice() const template void VectorHIP::resize(const int size) { - Utils::HostHIP::gpuMemcpyHostToDevice(mDeviceSize, &size, sizeof(int)); + utils::host_hip::gpuMemcpyHostToDevice(mDeviceSize, &size, sizeof(int)); } template @@ -201,20 +201,20 @@ void VectorHIP::reset(const T* const source, const int size, const int initia { if (size > mCapacity) { if (mArrayPointer != nullptr) { - Utils::HostHIP::gpuFree(mArrayPointer); + utils::host_hip::gpuFree(mArrayPointer); } - Utils::HostHIP::gpuMalloc(reinterpret_cast(&mArrayPointer), size * sizeof(T)); + utils::host_hip::gpuMalloc(reinterpret_cast(&mArrayPointer), size * sizeof(T)); mCapacity = size; } if (source != nullptr) { - Utils::HostHIP::gpuMemcpyHostToDevice(mArrayPointer, source, size * sizeof(T)); - Utils::HostHIP::gpuMemcpyHostToDevice(mDeviceSize, &size, sizeof(int)); + utils::host_hip::gpuMemcpyHostToDevice(mArrayPointer, source, size * sizeof(T)); + utils::host_hip::gpuMemcpyHostToDevice(mDeviceSize, &size, sizeof(int)); } else { - Utils::HostHIP::gpuMemcpyHostToDevice(mDeviceSize, &initialSize, sizeof(int)); + utils::host_hip::gpuMemcpyHostToDevice(mDeviceSize, &initialSize, sizeof(int)); } } @@ -227,7 +227,7 @@ void VectorHIP::copyIntoVector(std::vector& destinationVector, const int s try { hostPrimitivePointer = static_cast(malloc(size * sizeof(T))); - Utils::HostHIP::gpuMemcpyDeviceToHost(hostPrimitivePointer, mArrayPointer, size * sizeof(T)); + utils::host_hip::gpuMemcpyDeviceToHost(hostPrimitivePointer, mArrayPointer, size * sizeof(T)); destinationVector = std::move(std::vector(hostPrimitivePointer, hostPrimitivePointer + size)); @@ -245,7 +245,7 @@ void VectorHIP::copyIntoVector(std::vector& destinationVector, const int s template void VectorHIP::copyIntoSizedVector(std::vector& destinationVector) { - Utils::HostHIP::gpuMemcpyDeviceToHost(destinationVector.data(), mArrayPointer, destinationVector.size() * sizeof(T)); + utils::host_hip::gpuMemcpyDeviceToHost(destinationVector.data(), mArrayPointer, destinationVector.size() * sizeof(T)); } template @@ -253,12 +253,12 @@ inline void VectorHIP::destroy() { if (mArrayPointer != nullptr) { - Utils::HostHIP::gpuFree(mArrayPointer); + utils::host_hip::gpuFree(mArrayPointer); } if (mDeviceSize != nullptr) { - Utils::HostHIP::gpuFree(mDeviceSize); + utils::host_hip::gpuFree(mDeviceSize); } } @@ -290,7 +290,7 @@ template T VectorHIP::getElementFromDevice(const int index) const { T element; - Utils::HostHIP::gpuMemcpyDeviceToHost(&element, mArrayPointer + index, sizeof(T)); + utils::host_hip::gpuMemcpyDeviceToHost(&element, mArrayPointer + index, sizeof(T)); return element; } @@ -304,7 +304,7 @@ GPUhd() int VectorHIP::size() const template GPUd() int VectorHIP::extend(const int sizeIncrement) const { - const int startIndex = Utils::DeviceHIP::gpuAtomicAdd(mDeviceSize, sizeIncrement); + const int startIndex = utils::device_hip::gpuAtomicAdd(mDeviceSize, sizeIncrement); assert(size() <= mCapacity); return startIndex; diff --git a/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/VertexerTraitsHIP.h b/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/VertexerTraitsHIP.h index 245af2485a55f..abeddc6594898 100644 --- a/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/VertexerTraitsHIP.h +++ b/Detectors/ITSMFT/ITS/tracking/hip/include/ITStrackingHIP/VertexerTraitsHIP.h @@ -21,7 +21,6 @@ #include "ITStracking/VertexerTraits.h" #include "ITStracking/Cluster.h" #include "ITStracking/Constants.h" -// #include "ITStracking/Definitions.h" #include "ITStracking/Tracklet.h" #include "ITStrackingHIP/DeviceStoreVertexerHIP.h" diff --git a/Detectors/ITSMFT/ITS/tracking/hip/src/ContextHIP.hip.cxx b/Detectors/ITSMFT/ITS/tracking/hip/src/ContextHIP.hip.cxx index dc6ea9cb3a08b..d43cfc2320ff2 100644 --- a/Detectors/ITSMFT/ITS/tracking/hip/src/ContextHIP.hip.cxx +++ b/Detectors/ITSMFT/ITS/tracking/hip/src/ContextHIP.hip.cxx @@ -43,7 +43,7 @@ namespace its namespace gpu { -using Utils::HostHIP::checkHIPError; +using utils::host_hip::checkHIPError; ContextHIP::ContextHIP(bool dumpDevices) { diff --git a/Detectors/ITSMFT/ITS/tracking/hip/src/DeviceStoreVertexerHIP.hip.cxx b/Detectors/ITSMFT/ITS/tracking/hip/src/DeviceStoreVertexerHIP.hip.cxx index de82d3835f7be..e3c00813ece87 100644 --- a/Detectors/ITSMFT/ITS/tracking/hip/src/DeviceStoreVertexerHIP.hip.cxx +++ b/Detectors/ITSMFT/ITS/tracking/hip/src/DeviceStoreVertexerHIP.hip.cxx @@ -84,8 +84,8 @@ UniquePointer DeviceStoreVertexerHIP::initialise(const s mIndexTables[0].reset(indexTables[0].data(), static_cast(indexTables[0].size())); mIndexTables[1].reset(indexTables[2].data(), static_cast(indexTables[2].size())); - const dim3 threadsPerBlock{Utils::HostHIP::getBlockSize(mClusters[1].capacity())}; - const dim3 blocksGrid{Utils::HostHIP::getBlocksGrid(threadsPerBlock, mClusters[1].capacity())}; + const dim3 threadsPerBlock{utils::host_hip::getBlockSize(mClusters[1].capacity())}; + const dim3 blocksGrid{utils::host_hip::getBlocksGrid(threadsPerBlock, mClusters[1].capacity())}; UniquePointer deviceStoreVertexerPtr{*this}; diff --git a/Detectors/ITSMFT/ITS/tracking/hip/src/UtilsHIP.hip.cxx b/Detectors/ITSMFT/ITS/tracking/hip/src/UtilsHIP.hip.cxx index da9a149181a60..fb568e3dfd773 100644 --- a/Detectors/ITSMFT/ITS/tracking/hip/src/UtilsHIP.hip.cxx +++ b/Detectors/ITSMFT/ITS/tracking/hip/src/UtilsHIP.hip.cxx @@ -8,7 +8,7 @@ // granted to it by virtue of its status as an Intergovernmental Organization // or submit itself to any jurisdiction. /// -/// \file UtilsHIP.hip.cxx +/// \file utilsHIP.hip.cxx /// \brief /// @@ -58,7 +58,7 @@ namespace its namespace gpu { -void Utils::HostHIP::checkHIPError(const hipError_t error, const char* file, const int line) +void utils::host_hip::checkHIPError(const hipError_t error, const char* file, const int line) { if (error != hipSuccess) { std::ostringstream errorString{}; @@ -68,18 +68,18 @@ void Utils::HostHIP::checkHIPError(const hipError_t error, const char* file, con } } -dim3 Utils::HostHIP::getBlockSize(const int colsNum) +dim3 utils::host_hip::getBlockSize(const int colsNum) { return getBlockSize(colsNum, 1); } -dim3 Utils::HostHIP::getBlockSize(const int colsNum, const int rowsNum) +dim3 utils::host_hip::getBlockSize(const int colsNum, const int rowsNum) { const DeviceProperties& deviceProperties = ContextHIP::getInstance().getDeviceProperties(); return getBlockSize(colsNum, rowsNum, deviceProperties.streamProcessors / deviceProperties.maxBlocksPerSM); } -dim3 Utils::HostHIP::getBlockSize(const int colsNum, const int rowsNum, const int maxThreadsPerBlock) +dim3 utils::host_hip::getBlockSize(const int colsNum, const int rowsNum, const int maxThreadsPerBlock) { const DeviceProperties& deviceProperties = ContextHIP::getInstance().getDeviceProperties(); int xThreads = std::max(std::min(colsNum, static_cast(deviceProperties.maxThreadsDim.x)), 1); @@ -98,57 +98,57 @@ dim3 Utils::HostHIP::getBlockSize(const int colsNum, const int rowsNum, const in return dim3{static_cast(xThreads), static_cast(yThreads)}; } -dim3 Utils::HostHIP::getBlocksGrid(const dim3& threadsPerBlock, const int rowsNum) +dim3 utils::host_hip::getBlocksGrid(const dim3& threadsPerBlock, const int rowsNum) { return getBlocksGrid(threadsPerBlock, rowsNum, 1); } -dim3 Utils::HostHIP::getBlocksGrid(const dim3& threadsPerBlock, const int rowsNum, const int colsNum) +dim3 utils::host_hip::getBlocksGrid(const dim3& threadsPerBlock, const int rowsNum, const int colsNum) { return dim3{1 + (rowsNum - 1) / threadsPerBlock.x, 1 + (colsNum - 1) / threadsPerBlock.y}; } -void Utils::HostHIP::gpuMalloc(void** p, const int size) +void utils::host_hip::gpuMalloc(void** p, const int size) { checkHIPError(hipMalloc(p, size), __FILE__, __LINE__); } -void Utils::HostHIP::gpuFree(void* p) +void utils::host_hip::gpuFree(void* p) { checkHIPError(hipFree(p), __FILE__, __LINE__); } -void Utils::HostHIP::gpuMemset(void* p, int value, int size) +void utils::host_hip::gpuMemset(void* p, int value, int size) { checkHIPError(hipMemset(p, value, size), __FILE__, __LINE__); } -void Utils::HostHIP::gpuMemcpyHostToDevice(void* dst, const void* src, int size) +void utils::host_hip::gpuMemcpyHostToDevice(void* dst, const void* src, int size) { checkHIPError(hipMemcpy(dst, src, size, hipMemcpyHostToDevice), __FILE__, __LINE__); } -void Utils::HostHIP::gpuMemcpyHostToDeviceAsync(void* dst, const void* src, int size, hipStream_t& stream) +void utils::host_hip::gpuMemcpyHostToDeviceAsync(void* dst, const void* src, int size, hipStream_t& stream) { checkHIPError(hipMemcpyAsync(dst, src, size, hipMemcpyHostToDevice, stream), __FILE__, __LINE__); } -void Utils::HostHIP::gpuMemcpyDeviceToHost(void* dst, const void* src, int size) +void utils::host_hip::gpuMemcpyDeviceToHost(void* dst, const void* src, int size) { checkHIPError(hipMemcpy(dst, src, size, hipMemcpyDeviceToHost), __FILE__, __LINE__); } -// void Utils::HostHIP::gpuStartProfiler() +// void utils::host_hip::gpuStartProfiler() // { // checkHIPError(hipProfilerStart(), __FILE__, __LINE__); // } -// void Utils::HostHIP::gpuStopProfiler() +// void utils::host_hip::gpuStopProfiler() // { // checkHIPError(hipProfilerStop(), __FILE__, __LINE__); // } -GPUd() int Utils::DeviceHIP::getLaneIndex() +GPUd() int utils::device_hip::getLaneIndex() { uint32_t laneIndex; asm volatile("mov.u32 %0, %%laneid;" @@ -156,13 +156,13 @@ GPUd() int Utils::DeviceHIP::getLaneIndex() return static_cast(laneIndex); } -// GPUd() int Utils::Device::shareToWarp(const int value, const int laneIndex) +// GPUd() int utils::Device::shareToWarp(const int value, const int laneIndex) // { // cooperative_groups::coalesced_group threadGroup = cooperative_groups::coalesced_threads(); // return threadGroup.shfl(value, laneIndex); // } -// GPUd() int Utils::Device::gpuAtomicAdd(int* p, const int incrementSize) +// GPUd() int utils::Device::gpuAtomicAdd(int* p, const int incrementSize) // { // return atomicAdd(p, incrementSize); // } diff --git a/Detectors/ITSMFT/ITS/tracking/hip/src/VertexerTraitsHIP.hip.cxx b/Detectors/ITSMFT/ITS/tracking/hip/src/VertexerTraitsHIP.hip.cxx index 1fc760a87d639..2b107653691d5 100644 --- a/Detectors/ITSMFT/ITS/tracking/hip/src/VertexerTraitsHIP.hip.cxx +++ b/Detectors/ITSMFT/ITS/tracking/hip/src/VertexerTraitsHIP.hip.cxx @@ -344,8 +344,8 @@ void VertexerTraitsHIP::computeTracklets() std::cout << "\t\tno clusters on layer 1. Returning.\n"; return; } - const dim3 threadsPerBlock{gpu::Utils::HostHIP::getBlockSize(mClusters[1].capacity())}; - const dim3 blocksGrid{gpu::Utils::HostHIP::getBlocksGrid(threadsPerBlock, mClusters[1].capacity())}; + const dim3 threadsPerBlock{gpu::utils::host_hip::getBlockSize(mClusters[1].capacity())}; + const dim3 blocksGrid{gpu::utils::host_hip::getBlocksGrid(threadsPerBlock, mClusters[1].capacity())}; hipLaunchKernelGGL((gpu::trackleterKernel), dim3(blocksGrid), dim3(threadsPerBlock), 0, 0, getDeviceContextPtr(), @@ -375,8 +375,8 @@ void VertexerTraitsHIP::computeTrackletMatching() std::cout << "\t\tno clusters on layer 1. Returning.\n"; return; } - const dim3 threadsPerBlock{gpu::Utils::HostHIP::getBlockSize(mClusters[1].capacity())}; - const dim3 blocksGrid{gpu::Utils::HostHIP::getBlocksGrid(threadsPerBlock, mClusters[1].capacity())}; + const dim3 threadsPerBlock{gpu::utils::host_hip::getBlockSize(mClusters[1].capacity())}; + const dim3 blocksGrid{gpu::utils::host_hip::getBlocksGrid(threadsPerBlock, mClusters[1].capacity())}; size_t bufferSize = mStoreVertexerGPU.getConfig().tmpCUBBufferSize * sizeof(int); hipLaunchKernelGGL((gpu::trackletSelectionKernel), dim3(blocksGrid), dim3(threadsPerBlock), 0, 0, @@ -423,8 +423,8 @@ void VertexerTraitsHIP::computeVertices() std::cout << "\t\tno clusters on layer 1. Returning.\n"; return; } - const dim3 threadsPerBlock{gpu::Utils::HostHIP::getBlockSize(mClusters[1].capacity())}; - const dim3 blocksGrid{gpu::Utils::HostHIP::getBlocksGrid(threadsPerBlock, mClusters[1].capacity())}; + const dim3 threadsPerBlock{gpu::utils::host_hip::getBlockSize(mClusters[1].capacity())}; + const dim3 blocksGrid{gpu::utils::host_hip::getBlocksGrid(threadsPerBlock, mClusters[1].capacity())}; size_t bufferSize = mStoreVertexerGPU.getConfig().tmpCUBBufferSize * sizeof(int); int nLines = mStoreVertexerGPU.getNExclusiveFoundLines().getElementFromDevice(mClusters[1].size() - 1) + mStoreVertexerGPU.getNFoundLines().getElementFromDevice(mClusters[1].size() - 1); int nCentroids{static_cast(nLines * (nLines - 1) / 2)};