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

[Backport 12_0_X] Fix for SiPixelRecHitFromCUDA crash during online GPU tests #35317

Merged
merged 1 commit into from Sep 20, 2021
Merged

[Backport 12_0_X] Fix for SiPixelRecHitFromCUDA crash during online GPU tests #35317

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Jump to
Jump to file
Failed to load files.
Diff view
Diff view
1 change: 1 addition & 0 deletions CUDADataFormats/SiPixelCluster/interface/gpuClusteringConstants.h
View file
Open in desktop
Expand Up @@ -18,6 +18,7 @@ namespace gpuClustering {
constexpr uint16_t maxNumModules = 2000;
constexpr int32_t maxNumClustersPerModules = maxHitsInModule();
constexpr uint16_t invalidModuleId = std::numeric_limits<uint16_t>::max() - 1;
constexpr int invalidClusterId = -9999;
static_assert(invalidModuleId > maxNumModules); // invalidModuleId must be > maxNumModules

} // namespace gpuClustering
Expand Down
14 changes: 12 additions & 2 deletions RecoLocalTracker/SiPixelClusterizer/plugins/SiPixelDigisClustersFromSoA.cc
View file
Open in desktop
Expand Up @@ -81,8 +81,14 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con
edm::DetSet<PixelDigi>* detDigis = nullptr;
uint32_t detId = 0;
for (uint32_t i = 0; i < nDigis; i++) {
if (digis.pdigi(i) == 0)
// check for uninitialized digis
// this is set in RawToDigi_kernel in SiPixelRawToClusterGPUKernel.cu
if (digis.rawIdArr(i) == 0)
continue;
// check for noisy/dead pixels (electrons set to 0)
if (digis.adc(i) == 0)
continue;

detId = digis.rawIdArr(i);
if (storeDigis_) {
detDigis = &collection->find_or_insert(detId);
Expand Down Expand Up @@ -134,7 +140,11 @@ void SiPixelDigisClustersFromSoA::produce(edm::StreamID, edm::Event& iEvent, con
};

for (uint32_t i = 0; i < nDigis; i++) {
if (digis.pdigi(i) == 0)
// check for uninitialized digis
if (digis.rawIdArr(i) == 0)
continue;
// check for noisy/dead pixels (electrons set to 0)
if (digis.adc(i) == 0)
continue;
if (digis.clus(i) > 9000)
continue; // not in cluster; TODO add an assert for the size
Expand Down
2 changes: 1 addition & 1 deletion RecoLocalTracker/SiPixelClusterizer/plugins/gpuCalibPixel.h
View file
Open in desktop
Expand Up @@ -54,9 +54,9 @@ namespace gpuCalibPixel {
float gain = ret.second;
// float pedestal = 0; float gain = 1.;
if (isDeadColumn | isNoisyColumn) {
printf("bad pixel at %d in %d\n", i, id[i]);
id[i] = invalidModuleId;
adc[i] = 0;
printf("bad pixel at %d in %d\n", i, id[i]);
} else {
float vcal = adc[i] * gain - pedestal * gain;
adc[i] = std::max(100, int(vcal * conversionFactor + offset));
Expand Down
2 changes: 1 addition & 1 deletion RecoLocalTracker/SiPixelClusterizer/plugins/gpuClustering.h
View file
Open in desktop
Expand Up @@ -275,7 +275,7 @@ namespace gpuClustering {
// adjust the cluster id to be a positive value starting from 0
for (int i = first; i < msize; i += blockDim.x) {
if (id[i] == invalidModuleId) { // skip invalid pixels
clusterId[i] = -9999;
clusterId[i] = invalidClusterId;
continue;
}
clusterId[i] = -clusterId[i] - 1;
Expand Down