From 30f4f4b81e0f9a5087c4aafeb2158c04af39799c Mon Sep 17 00:00:00 2001
From: dsperka <dsperka@cern.ch>
Date: Tue, 16 Apr 2019 11:18:53 -0400
Subject: [PATCH] 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).
---
 .../plugins/PixelTrackReconstructionGPU.cc             |  4 ++--
 .../plugins/PixelTrackReconstructionGPU.cu             | 10 +++++-----
 2 files changed, 7 insertions(+), 7 deletions(-)

diff --git a/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackReconstructionGPU.cc b/RecoPixelVertexing/PixelTrackFitting/plugins/PixelTrackReconstructionGPU.cc
index 2a360b1da58a6..33821851ace72 100644
--- 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
index 19a91ba8c83b4..2e3d296e0df6e 100644
--- 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>)));