diff --git a/Trigger/TrigAccel/TrigAccelEvent/TrigAccelEvent/TrigITkAccelEDM.h b/Trigger/TrigAccel/TrigAccelEvent/TrigAccelEvent/TrigITkAccelEDM.h
index a97a6a40bffea9378f1ba6fdd20d15dbff21b3f4..9b2e569e49a7c684ee0249339c2ab7b51a77cc53 100644
--- a/Trigger/TrigAccel/TrigAccelEvent/TrigAccelEvent/TrigITkAccelEDM.h
+++ b/Trigger/TrigAccel/TrigAccelEvent/TrigAccelEvent/TrigITkAccelEDM.h
@@ -18,7 +18,7 @@ namespace ITk {
static constexpr unsigned int MAX_NUMBER_SCT_MODULES = 24600;
static constexpr unsigned int MAX_NUMBER_SPACEPOINTS = 300000;
static constexpr unsigned int MAX_PHI_SLICES = 100;
- static constexpr unsigned int MAX_NUMBER_OUTPUT_SEEDS = 500000;
+ static constexpr unsigned int MAX_NUMBER_OUTPUT_SEEDS = 100000;
typedef struct SiliconLayer {
public:
@@ -92,7 +92,6 @@ namespace ITk {
int m_nSeeds;
int m_nMiddleSps;
int m_nI, m_nO;
- int m_nCandidates;
int m_nErrors;
int m_innerIndex[MAX_NUMBER_OUTPUT_SEEDS];
int m_middleIndex[MAX_NUMBER_OUTPUT_SEEDS];
diff --git a/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletCountingKernelCuda_ITk.cuh b/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletCountingKernelCuda_ITk.cuh
index 31347eed3bd7e84cde013d96e50ecb5d2092c8ab..c7fd6e4f81919f4446def34a828e9f61b99774dd 100644
--- a/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletCountingKernelCuda_ITk.cuh
+++ b/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletCountingKernelCuda_ITk.cuh
@@ -46,8 +46,10 @@ __global__ static void doubletCountingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SET
const float zPlus = dSettings->m_zedPlus + zTolerance;
const float maxTheta = 2*atan(exp(-maxEta));
const float maxCtg = cos(maxTheta)/sin(maxTheta);
- const float minOuterZ = dSettings->m_zedMinus - maxOuterRadius*maxCtg - zTolerance;
- const float maxOuterZ = dSettings->m_zedPlus + maxOuterRadius*maxCtg + zTolerance;
+ const float minZ0 = dSettings->m_zedMinus;
+ const float maxZ0 = dSettings->m_zedPlus;
+ const float minOuterZ = minZ0 - maxOuterRadius*maxCtg - zTolerance;
+ const float maxOuterZ = maxZ0 + maxOuterRadius*maxCtg + zTolerance;
const double ptCoeff = 0.29997*dSettings->m_magFieldZ/2.0;// ~0.3*B/2 - assumes nominal field of 2*T
const float tripletPtMin = dSettings->m_tripletPtMin; // Retrieve from settings
@@ -62,6 +64,8 @@ __global__ static void doubletCountingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SET
float zm = dSpacepoints->m_z[spmIdx];
float rm = dSpacepoints->m_r[spmIdx];
+ if (!canBeMiddleSpacePoint(rm, zm)) continue;
+
if(threadIdx.y ==0) {
nInner[threadIdx.x] = 0;
nOuter[threadIdx.x] = 0;
@@ -88,11 +92,10 @@ __global__ static void doubletCountingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SET
if(next_spEnd == next_spBegin) continue;//no spacepoints in this layer
const TrigAccel::ITk::SILICON_LAYER& layerGeo = dDetModel->m_layers[nextLayerIdx];
+
bool isBarrel = (layerGeo.m_type == 0);
-
float refCoord = layerGeo.m_refCoord;
-
if(isBarrel && std::abs(refCoord-rm)>maxDoubletLength) continue;
//boundaries for nextLayer
@@ -139,6 +142,9 @@ __global__ static void doubletCountingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SET
if(std::abs(tau)>maxCtg) continue;
// Cut on Z
+ float z0 = zsp - rsp*tau;
+ if(z0 < minZ0 || z0 > maxZ0) continue;
+
float outZ = zsp + (maxOuterRadius-rsp)*tau;
if(outZ < minOuterZ || outZ > maxOuterZ) continue;
diff --git a/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletHelperFunctionsCuda_ITk.cuh b/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletHelperFunctionsCuda_ITk.cuh
index e8cc2308d67714a2faac4261018360d259e8b0d0..25e8ca00f21c8735b1889a00fac5bdd95501e078 100644
--- a/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletHelperFunctionsCuda_ITk.cuh
+++ b/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletHelperFunctionsCuda_ITk.cuh
@@ -30,4 +30,18 @@ __device__ static float getMaxDeltaLEta (float eta) {
else return eta*eta*eta*eta*1.7582417 + eta*eta*-129.67033 + 3324.61538;
}
+
+__device__ static float canBeMiddleSpacePoint (float r, float z) {
+ float z_abs = std::abs(z);
+
+ // Barrel
+ if ((r < 90.) && (z_abs < 245.)) return false; // volume 80
+ if ((r > 200.) && (z_abs < 245.)) return false; // volume 83
+ if ((r > 250.) && (z_abs < 350.)) return false; // volume 84
+ if (r > 290.) return false; // volume 84
+ // End Cap
+ if (z_abs > 2600.) return false;
+ return true;
+}
+
#endif
\ No newline at end of file
diff --git a/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletMakingKernelCuda_ITk.cuh b/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletMakingKernelCuda_ITk.cuh
index 715d0ad8c60a68d306d7d45c7d7c312001669aa3..c0a3fdfb9ba1083daabab975cded93a0f4d6caf8 100644
--- a/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletMakingKernelCuda_ITk.cuh
+++ b/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletMakingKernelCuda_ITk.cuh
@@ -49,8 +49,10 @@ __global__ static void doubletMakingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SETTI
const float zPlus = dSettings->m_zedPlus + zTolerance;
const float maxTheta = 2*atan(exp(-maxEta));
const float maxCtg = cos(maxTheta)/sin(maxTheta);
- const float minOuterZ = dSettings->m_zedMinus - maxOuterRadius*maxCtg - zTolerance;
- const float maxOuterZ = dSettings->m_zedPlus + maxOuterRadius*maxCtg + zTolerance;
+ const float minZ0 = dSettings->m_zedMinus;
+ const float maxZ0 = dSettings->m_zedPlus;
+ const float minOuterZ = minZ0 - maxOuterRadius*maxCtg - zTolerance;
+ const float maxOuterZ = maxZ0 + maxOuterRadius*maxCtg + zTolerance;
const double ptCoeff = 0.29997*dSettings->m_magFieldZ/2.0;// ~0.3*B/2 - assumes nominal field of 2*T
const float tripletPtMin = dSettings->m_tripletPtMin; // Retrieve from settings
@@ -89,6 +91,8 @@ __global__ static void doubletMakingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SETTI
float zm = dSpacepoints->m_z[spmIdx];
float rm = dSpacepoints->m_r[spmIdx];
+ if (!canBeMiddleSpacePoint(rm, zm)) continue;
+
//2. loop over other phi-bins / layers
for(int deltaPhiIdx=-1;deltaPhiIdx<=1;deltaPhiIdx++) {
@@ -108,8 +112,8 @@ __global__ static void doubletMakingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SETTI
const TrigAccel::ITk::SILICON_LAYER& layerGeo = dDetModel->m_layers[nextLayerIdx];
bool isBarrel = (layerGeo.m_type == 0);
+
float refCoord = layerGeo.m_refCoord;
-
if(isBarrel && std::abs(refCoord-rm)>maxDoubletLength) continue;
//boundaries for nextLayer
@@ -156,6 +160,9 @@ __global__ static void doubletMakingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SETTI
if(std::abs(tau)>maxCtg) continue;
// Cut on Z
+ float z0 = zsp - rsp*tau;
+ if(z0 < minZ0 || z0 > maxZ0) continue;
+
float outZ = zsp + (maxOuterRadius-rsp)*tau;
if(outZ < minOuterZ || outZ > maxOuterZ) continue;
diff --git a/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletMatchingKernelCuda_ITk.cuh b/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletMatchingKernelCuda_ITk.cuh
index 8c803fd3068ba66467736a114d7d97cc00810db9..68fc24c66a10c77c54841c2c3e3d4030ab995ad8 100644
--- a/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletMatchingKernelCuda_ITk.cuh
+++ b/Trigger/TrigAccel/TrigInDetCUDA/src/DoubletMatchingKernelCuda_ITk.cuh
@@ -8,6 +8,7 @@
#include <cuda_runtime.h>
#include "SeedMakingDataStructures_ITk.h"
+#include "DoubletHelperFunctionsCuda_ITk.cuh"
__global__ static void doubletMatchingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SETTINGS* dSettings,
@@ -25,6 +26,7 @@ __global__ static void doubletMatchingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SET
__shared__ int spmIdx;
__shared__ float rm;
+ __shared__ float zm;
__shared__ float covZ;
__shared__ float covR;
@@ -91,6 +93,7 @@ __global__ static void doubletMatchingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SET
outerStart = d_Storage->m_outerStart[itemIdx];
rm = dSpacepoints->m_r[spmIdx];
+ zm = dSpacepoints->m_z[spmIdx];
covZ = dSpacepoints->m_covZ[spmIdx];
covR = dSpacepoints->m_covR[spmIdx];
@@ -100,7 +103,9 @@ __global__ static void doubletMatchingKernel_ITk(TrigAccel::ITk::SEED_FINDER_SET
cosA = x0/rm;
sinA = y0/rm;
}
- __syncthreads();
+ __syncthreads();
+
+ if (!canBeMiddleSpacePoint(rm, zm)) continue;
for(int innerIdx = threadIdx.x; innerIdx<nInner;innerIdx+=blockDim.x) {
diff --git a/Trigger/TrigAccel/TrigInDetCUDA/src/SeedMakingDataStructures_ITk.h b/Trigger/TrigAccel/TrigInDetCUDA/src/SeedMakingDataStructures_ITk.h
index 57520fcb5e88cc6e5dc127162b084968572eb8ef..94e88de1492bf0bc27a986c8c94e17ae7f9101c3 100644
--- a/Trigger/TrigAccel/TrigInDetCUDA/src/SeedMakingDataStructures_ITk.h
+++ b/Trigger/TrigAccel/TrigInDetCUDA/src/SeedMakingDataStructures_ITk.h
@@ -8,7 +8,7 @@
#include "CommonStructures.h"
static constexpr unsigned int MAX_MIDDLE_SP_ITk = 300000;
-static constexpr unsigned int MAX_DOUBLET_ITk = 50000000;
+static constexpr unsigned int MAX_DOUBLET_ITk = 25000000;
static constexpr unsigned int NUM_MIDDLE_THREADS_ITk = 32;
static constexpr unsigned int OUTER_THREADS_MULTIPLIER_ITk = 4; // i.e thread block is 32 x 4*192/32
static constexpr unsigned int MAX_NUMBER_DOUBLETS_ITk = 1500;
diff --git a/Trigger/TrigAccel/TrigInDetCUDA/src/SeedMakingWorkCuda_ITk.cu b/Trigger/TrigAccel/TrigInDetCUDA/src/SeedMakingWorkCuda_ITk.cu
index f654ec73e45311babb8826753d1ed2b09ce23f3d..03645ebb4ce3ebef922de88a2740d005433e045d 100644
--- a/Trigger/TrigAccel/TrigInDetCUDA/src/SeedMakingWorkCuda_ITk.cu
+++ b/Trigger/TrigAccel/TrigInDetCUDA/src/SeedMakingWorkCuda_ITk.cu
@@ -151,7 +151,6 @@ bool SeedMakingWorkCudaITk::run() {
pOutput->m_nSeeds = 0;
pOutput->m_nI = 0;
pOutput->m_nO = 0;
- pOutput->m_nCandidates = 0;
cudaMemcpyAsync(p.h_outputseeds, p.d_outputseeds, sizeof(TrigAccel::ITk::OUTPUT_SEED_STORAGE), cudaMemcpyDeviceToHost, p.m_stream);
@@ -297,7 +296,6 @@ bool SeedMakingWorkCudaManagedITk::run() {
pOutput->m_nSeeds = 0;
pOutput->m_nI = 0;
pOutput->m_nO = 0;
- pOutput->m_nCandidates = 0;
checkError();