Skip to content
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

Jump to bottom

move to int32 for hit indices #583

Merged
merged 5 commits into from
Nov 27, 2020
Merged

move to int32 for hit indices #583

Show file tree
Hide file tree
Changes from 2 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
23c546e
increase maxhit, count forlon doublets
VinInn Jun 6, 2020
7877d3a
fix hit index in phi binner
VinInn Jun 7, 2020
1af9050
Remove commented out code
fwyzard Nov 27, 2020
62a0c30
Apply code formatting
fwyzard Nov 27, 2020
2c8573d
Apply code formatting
fwyzard Nov 27, 2020
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
2 changes: 1 addition & 1 deletion CUDADataFormats/Track/interface/PixelTrackHeterogeneous.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,7 @@ class TrackSoAT {
static constexpr int32_t stride() { return S; }

using Quality = trackQuality::Quality;
using hindex_type = uint16_t;
using hindex_type = uint32_t;
using HitContainer = cms::cuda::OneToManyAssoc<hindex_type, S, 5 * S>;

// Always check quality is at least loose!
Expand Down
10 changes: 6 additions & 4 deletions CUDADataFormats/TrackingRecHit/interface/TrackingRecHit2DSOAView.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -15,10 +15,12 @@ namespace pixelCPEforGPU {
class TrackingRecHit2DSOAView {
public:
static constexpr uint32_t maxHits() { return gpuClustering::MaxNumClusters; }
using hindex_type = uint16_t; // if above is <=2^16
using hindex_type = uint32_t; // if above is <=2^32

using Hist =
cms::cuda::HistoContainer<int16_t, 128, gpuClustering::MaxNumClusters, 8 * sizeof(int16_t), uint16_t, 10>;
using PhiBinner =
cms::cuda::HistoContainer<int16_t, 128, gpuClustering::MaxNumClusters, 8 * sizeof(int16_t), hindex_type, 10>;

using Hist = PhiBinner; // FIXME

using AverageGeometry = phase1PixelTopology::AverageGeometry;

Expand Down Expand Up @@ -93,7 +95,7 @@ class TrackingRecHit2DSOAView {

uint32_t* m_hitsLayerStart;

Hist* m_hist;
PhiBinner * m_hist; // FIXME use a more descriptive name consistently

uint32_t m_nHits;
};
Expand Down
2 changes: 1 addition & 1 deletion RecoPixelVertexing/PixelTriplets/plugins/CAConstants.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -44,7 +44,7 @@ namespace CAConstants {
constexpr uint32_t maxTuples() { return maxNumberOfTuples(); }

// types
using hindex_type = uint16_t; // FIXME from siPixelRecHitsHeterogeneousProduct
using hindex_type = uint32_t; // FIXME from siPixelRecHitsHeterogeneousProduct
using tindex_type = uint16_t; // for tuples

#ifndef ONLY_PHICUT
Expand Down
1 change: 1 addition & 0 deletions RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -134,6 +134,7 @@ void CAHitNtupletGeneratorKernelsCPU::launchKernels(HitsOnCPU const &hh, TkSoA *
if (m_params.doStats_) {
kernel_checkOverflows(tuples_d,
device_tupleMultiplicity_.get(),
device_hitToTuple_.get(),
device_hitTuple_apc_,
device_theCells_.get(),
device_nCells_,
Expand Down
33 changes: 18 additions & 15 deletions RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernels.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -125,21 +125,6 @@ void CAHitNtupletGeneratorKernelsGPU::launchKernels(HitsOnCPU const &hh, TkSoA *
cudaCheck(cudaGetLastError());
}

if (m_params.doStats_) {
numberOfBlocks = (std::max(nhits, m_params.maxNumberOfDoublets_) + blockSize - 1) / blockSize;
kernel_checkOverflows<<<numberOfBlocks, blockSize, 0, cudaStream>>>(tuples_d,
device_tupleMultiplicity_.get(),
device_hitTuple_apc_,
device_theCells_.get(),
device_nCells_,
device_theCellNeighbors_.get(),
device_theCellTracks_.get(),
device_isOuterHitOfCell_.get(),
nhits,
m_params.maxNumberOfDoublets_,
counters_);
cudaCheck(cudaGetLastError());
}
#ifdef GPU_DEBUG
cudaDeviceSynchronize();
cudaCheck(cudaGetLastError());
Expand Down Expand Up @@ -278,6 +263,24 @@ void CAHitNtupletGeneratorKernelsGPU::classifyTuples(HitsOnCPU const &hh, TkSoA
cudaCheck(cudaGetLastError());
}

if (m_params.doStats_) {
auto nhits = hh.nHits();
numberOfBlocks = (std::max(nhits, m_params.maxNumberOfDoublets_) + blockSize - 1) / blockSize;
kernel_checkOverflows<<<numberOfBlocks, blockSize, 0, cudaStream>>>(tuples_d,
device_tupleMultiplicity_.get(),
device_hitToTuple_.get(),
device_hitTuple_apc_,
device_theCells_.get(),
device_nCells_,
device_theCellNeighbors_.get(),
device_theCellTracks_.get(),
device_isOuterHitOfCell_.get(),
nhits,
m_params.maxNumberOfDoublets_,
counters_);
cudaCheck(cudaGetLastError());
}

if (m_params.doStats_) {
// counters (add flag???)
numberOfBlocks = (HitToTuple::capacity() + blockSize - 1) / blockSize;
Expand Down
9 changes: 7 additions & 2 deletions RecoPixelVertexing/PixelTriplets/plugins/CAHitNtupletGeneratorKernelsImpl.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -30,7 +30,8 @@ using TkSoA = pixelTrack::TrackSoA;
using HitContainer = pixelTrack::HitContainer;

__global__ void kernel_checkOverflows(HitContainer const *foundNtuplets,
CAConstants::TupleMultiplicity *tupleMultiplicity,
CAConstants::TupleMultiplicity const * tupleMultiplicity,
CAHitNtupletGeneratorKernelsGPU::HitToTuple const * hitToTuple,
cms::cuda::AtomicPairCounter *apc,
GPUCACell const *__restrict__ cells,
uint32_t const *__restrict__ nCells,
Expand Down Expand Up @@ -95,7 +96,11 @@ __global__ void kernel_checkOverflows(HitContainer const *foundNtuplets,
atomicAdd(&c.nKilledCells, 1);
if (0 == thisCell.theUsed)
atomicAdd(&c.nEmptyCells, 1);
if (thisCell.tracks().empty())
// if (thisCell.tracks().empty())
// atomicAdd(&c.nZeroTrackCells, 1);
fwyzard marked this conversation as resolved.
Show resolved Hide resolved
if (0==hitToTuple->size(thisCell.get_inner_hit_id()) &&
0==hitToTuple->size(thisCell.get_outer_hit_id())
)
atomicAdd(&c.nZeroTrackCells, 1);
}

Expand Down