diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc index 93b92e145ec5c..5a125068f9a85 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterCUDA.cc @@ -35,6 +35,7 @@ #include "HeterogeneousCore/CUDACore/interface/ScopedContext.h" #include "HeterogeneousCore/CUDAServices/interface/CUDAService.h" #include "RecoTracker/Record/interface/CkfComponentsRecord.h" +#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h" // local includes #include "SiPixelRawToClusterGPUKernel.h" @@ -77,6 +78,8 @@ class SiPixelRawToClusterCUDA : public edm::stream::EDProducer("CablingMapLabel")))), isRun2_(iConfig.getParameter("isRun2")), includeErrors_(iConfig.getParameter("IncludeErrors")), - useQuality_(iConfig.getParameter("UseQualityInfo")) { + useQuality_(iConfig.getParameter("UseQualityInfo")), + clusterThresholds_(iConfig.getParameter("theGPUClusterThreshold_L1"), + iConfig.getParameter("theGPUClusterThreshold")) { if (includeErrors_) { digiErrorPutToken_ = produces>(); } @@ -110,6 +115,8 @@ void SiPixelRawToClusterCUDA::fillDescriptions(edm::ConfigurationDescriptions& d desc.add("isRun2", true); desc.add("IncludeErrors", true); desc.add("UseQualityInfo", false); + desc.add("theGPUClusterThreshold_L1", 2000); + desc.add("theGPUClusterThreshold", 4000); desc.add("InputLabel", edm::InputTag("rawDataCollector")); { edm::ParameterSetDescription psd0; @@ -231,6 +238,7 @@ void SiPixelRawToClusterCUDA::acquire(const edm::Event& iEvent, } // end of for loop gpuAlgo_.makeClustersAsync(isRun2_, + clusterThresholds_, gpuMap, gpuModulesToUnpack, gpuGains, diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu index 25e5c925990f8..89b92af95398a 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.cu @@ -510,6 +510,7 @@ namespace pixelgpudetails { // Interface to outside void SiPixelRawToClusterGPUKernel::makeClustersAsync(bool isRun2, + const gpuClusterThresholds clusterThresholds, const SiPixelROCsStatusAndMapping *cablingMap, const unsigned char *modToUnp, const SiPixelGainForHLTonGPU *gains, @@ -635,7 +636,8 @@ namespace pixelgpudetails { #endif // apply charge cut - clusterChargeCut<<>>(digis_d.moduleInd(), + clusterChargeCut<<>>(clusterThresholds, + digis_d.moduleInd(), digis_d.adc(), clusters_d.moduleStart(), clusters_d.clusInModule(), diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h index 75eeab2606dd5..739c82ab0e17c 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelRawToClusterGPUKernel.h @@ -13,6 +13,7 @@ #include "HeterogeneousCore/CUDAUtilities/interface/host_noncached_unique_ptr.h" #include "DataFormats/SiPixelRawData/interface/SiPixelErrorCompact.h" #include "DataFormats/SiPixelRawData/interface/SiPixelFormatterErrors.h" +#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h" struct SiPixelROCsStatusAndMapping; class SiPixelGainForHLTonGPU; @@ -170,6 +171,7 @@ namespace pixelgpudetails { SiPixelRawToClusterGPUKernel& operator=(SiPixelRawToClusterGPUKernel&&) = delete; void makeClustersAsync(bool isRun2, + const gpuClusterThresholds clusterThresholds, const SiPixelROCsStatusAndMapping* cablingMap, const unsigned char* modToUnp, const SiPixelGainForHLTonGPU* gains, diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h index d9520da80b695..31f1b259fd9ba 100644 --- a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h @@ -7,12 +7,14 @@ #include "CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h" #include "HeterogeneousCore/CUDAUtilities/interface/cuda_assert.h" #include "HeterogeneousCore/CUDAUtilities/interface/prefixScan.h" +#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h" namespace gpuClustering { __global__ void clusterChargeCut( - uint16_t* __restrict__ id, // module id of each pixel (modified if bad cluster) - uint16_t const* __restrict__ adc, // charge of each pixel + gpuClusterThresholds clusterThresholds, // charge cut on cluster in electrons (for layer 1 and for other layers) + uint16_t* __restrict__ id, // module id of each pixel (modified if bad cluster) + uint16_t const* __restrict__ adc, // charge of each pixel uint32_t const* __restrict__ moduleStart, // index of the first pixel of each module uint32_t* __restrict__ nClustersInModule, // modified: number of clusters found in each module uint32_t const* __restrict__ moduleId, // module id of each module @@ -79,7 +81,7 @@ namespace gpuClustering { } __syncthreads(); - auto chargeCut = thisModuleId < 96 ? 2000 : 4000; // move in constants (calib?) + auto chargeCut = thisModuleId < 96 ? clusterThresholds.layer1 : clusterThresholds.otherLayers; for (auto i = threadIdx.x; i < nclus; i += blockDim.x) { newclusId[i] = ok[i] = charge[i] > chargeCut ? 1 : 0; } diff --git a/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h new file mode 100644 index 0000000000000..0222774993352 --- /dev/null +++ b/RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h @@ -0,0 +1,11 @@ +#ifndef RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterThresholds_h +#define RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterThresholds_h + +struct gpuClusterThresholds { + constexpr gpuClusterThresholds(int32_t clusterThreshold_L1, int32_t clusterThreshold): + layer1(clusterThreshold_L1), otherLayers(clusterThreshold){}; + const int32_t layer1; // Cluster threshold in electrons for Layer 1 + const int32_t otherLayers; // Cluster threshold in electrons +}; + +#endif // RecoLocalTracker_SiPixelClusterizer_plugins_gpuClusterThresholds_h \ No newline at end of file diff --git a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h index 02611ab1cac1d..b0d4b3bc4b693 100644 --- a/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h +++ b/RecoLocalTracker/SiPixelClusterizer/test/gpuClustering_t.h @@ -18,6 +18,7 @@ #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h" #include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterChargeCut.h" +#include "RecoLocalTracker/SiPixelClusterizer/plugins/gpuClusterThresholds.h" int main(void) { #ifdef __CUDACC__ @@ -28,6 +29,8 @@ int main(void) { constexpr int numElements = 256 * maxNumModules; + constexpr gpuClusterThresholds clusterThresholds(2000, 4000); + // these in reality are already on GPU auto h_id = std::make_unique(numElements); auto h_x = std::make_unique(numElements); @@ -288,6 +291,7 @@ int main(void) { cms::cuda::launch(clusterChargeCut, {blocksPerGrid, threadsPerBlock}, + clusterThresholds, d_id.get(), d_adc.get(), d_moduleStart.get(), @@ -317,8 +321,14 @@ int main(void) { if (ncl != std::accumulate(nclus, nclus + maxNumModules, 0)) std::cout << "ERROR!!!!! wrong number of cluster found" << std::endl; - clusterChargeCut( - h_id.get(), h_adc.get(), h_moduleStart.get(), h_clusInModule.get(), h_moduleId.get(), h_clus.get(), n); + clusterChargeCut(clusterThresholds, + h_id.get(), + h_adc.get(), + h_moduleStart.get(), + h_clusInModule.get(), + h_moduleId.get(), + h_clus.get(), + n); #endif // __CUDACC__ std::cout << "found " << nModules << " Modules active" << std::endl;