Fix the number of floats per hit (#321)

Fix the number of floats passed per hit (3 floats for GlobalPosition and 6 floats for GlobalError).
cms-patatrack · Apr 16, 2019 · 30f4f4b · 30f4f4b
1 parent 49f2c7f
commit 30f4f4b
Show file tree

Hide file tree

Showing 2 changed files with 7 additions and 7 deletions.
diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackReconstructionGPU.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackReconstructionGPU.cc
@@ -78,7 +78,7 @@ void PixelTrackReconstructionGPU::run(TracksWithTTRHs& tracks,
   // We use 3 floats for GlobalPosition and 6 floats for GlobalError (that's what is used by the Riemann fit).
   // Assume a safe maximum of 3K seeds: it will dynamically grow, if needed.
   int total_seeds = 0;
-  hits_and_covariances.reserve(sizeof(float)*3000*(points_in_seed*12));
+  hits_and_covariances.reserve(sizeof(float)*3000*(points_in_seed*9));
   for (auto const & regionHitSets : hitSets) {
     const TrackingRegion& region = regionHitSets.region();
     for (auto const & tuplet : regionHitSets) {
@@ -112,7 +112,7 @@ void PixelTrackReconstructionGPU::run(TracksWithTTRHs& tracks,
       sizeof(float)*hits_and_covariances.size(), cudaMemcpyDefault));
 
   // LAUNCH THE KERNEL FIT
-  launchKernelFit(hits_and_covariancesGPU, 12*4*total_seeds, 4,
+  launchKernelFit(hits_and_covariancesGPU, 9*4*total_seeds, 4,
       bField, helix_fit_resultsGPU);
   // CUDA MEMCOPY DEVICE2HOST OF HELIX_FIT
   cudaCheck(cudaDeviceSynchronize());

diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackReconstructionGPU.cu b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackReconstructionGPU.cu
@@ -25,7 +25,7 @@ KernelFastFitAllHits(float *hits_and_covariances,
   // Loop for hits_in_fit times:
   //   first 3 are the points
   //   the rest is the covariance matrix, 3x3
-  int start = (blockIdx.x * blockDim.x + threadIdx.x) * hits_in_fit * 12;
+  int start = (blockIdx.x * blockDim.x + threadIdx.x) * hits_in_fit * 9;
   int helix_start = (blockIdx.x * blockDim.x + threadIdx.x);
   if (start >= cumulative_size) {
     return;
@@ -67,7 +67,7 @@ KernelCircleFitAllHits(float *hits_and_covariances, int hits_in_fit,
   // Loop for hits_in_fit times:
   //   first 3 are the points
   //   the rest is the covariance matrix, 3x3
-  int start = (blockIdx.x * blockDim.x + threadIdx.x) * hits_in_fit * 12;
+  int start = (blockIdx.x * blockDim.x + threadIdx.x) * hits_in_fit * 9;
   int helix_start = (blockIdx.x * blockDim.x + threadIdx.x);
   if (start >= cumulative_size) {
     return;
@@ -115,7 +115,7 @@ KernelLineFitAllHits(float *hits_and_covariances, int hits_in_fit,
   // Loop for hits_in_fit times:
   //   first 3 are the points
   //   the rest is the covariance matrix, 3x3
-  int start = (blockIdx.x * blockDim.x + threadIdx.x) * hits_in_fit * 12;
+  int start = (blockIdx.x * blockDim.x + threadIdx.x) * hits_in_fit * 9;
   int helix_start = (blockIdx.x * blockDim.x + threadIdx.x);
   if (start >= cumulative_size) {
     return;
@@ -162,8 +162,8 @@ void PixelTrackReconstructionGPU::launchKernelFit(
     float B, Rfit::helix_fit *results)
 {
   const dim3 threads_per_block(32, 1);
-  int num_blocks = cumulative_size / (hits_in_fit * 12) / threads_per_block.x + 1;
-  auto numberOfSeeds = cumulative_size / (hits_in_fit * 12);
+  int num_blocks = cumulative_size / (hits_in_fit * 9) / threads_per_block.x + 1;
+  auto numberOfSeeds = cumulative_size / (hits_in_fit * 9);
 
   Rfit::Matrix3xNd<4> *hitsGPU;
   cudaCheck(cudaMalloc(&hitsGPU, 48 * numberOfSeeds * sizeof(Rfit::Matrix3xNd<4>)));