Skip to content

Migration of SiStripClusterizer to Alpaka #134

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

Open
mmasciov wants to merge 4 commits into
base: from-CMSSW_13_1_0
Choose a base branch
from
Open

Migration of SiStripClusterizer to Alpaka #134

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
a638875
First attempt to port to Alpaka SiStripClusters data format
richderem Aug 3, 2023
867b6fb
Renaming, cleaning and first working version
VourMa Aug 15, 2023
97d47db
Add missing using
VourMa Aug 15, 2023
22aa2a9
Further small cleanup of SiStripClustersX
mmasciov Aug 21, 2023
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions CUDADataFormats/SiStripCluster/BuildFile.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,7 @@
<use name="DataFormats/Common"/>
<use name="CUDADataFormats/Common"/>
<use name="HeterogeneousCore/CUDAUtilities"/>
<use name="DataFormats/SoATemplate" source_only="1"/>
<use name="cuda"/>
<use name="rootcore"/>

Expand Down
59 changes: 0 additions & 59 deletions CUDADataFormats/SiStripCluster/interface/SiStripClustersCUDA.h
View file
Open in desktop

This file was deleted.

35 changes: 35 additions & 0 deletions CUDADataFormats/SiStripCluster/interface/SiStripClustersSoADevice.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,35 @@
#ifndef CUDADataFormats_SiStripCluster_interface_SiStripClustersSoADevice_h
#define CUDADataFormats_SiStripCluster_interface_SiStripClustersSoADevice_h

#include "CUDADataFormats/Common/interface/PortableDeviceCollection.h"
#include "CUDADataFormats/SiStripCluster/interface/SiStripClustersSoAUtilities.h"

class SiStripClustersSoADevice : public cms::cuda::PortableDeviceCollection<SiStripClustersLayout> {
public:
using cms::cuda::PortableDeviceCollection<SiStripClustersLayout>::view;
using cms::cuda::PortableDeviceCollection<SiStripClustersLayout>::const_view;
using cms::cuda::PortableDeviceCollection<SiStripClustersLayout>::buffer;
using cms::cuda::PortableDeviceCollection<SiStripClustersLayout>::bufferSize;

SiStripClustersSoADevice() = default;
~SiStripClustersSoADevice() = default;

explicit SiStripClustersSoADevice(uint32_t maxClusters, cudaStream_t stream)
: cms::cuda::PortableDeviceCollection<SiStripClustersLayout>(maxClusters, stream), maxClusters_{maxClusters} {
const uint32_t maxStripsPerCluster = SiStripClustersSoA::maxStripsPerCluster; //768
cudaCheck(
cudaMemcpyAsync(&(view().maxClusterSize()), &maxStripsPerCluster, sizeof(uint32_t), cudaMemcpyDefault, stream));
};

SiStripClustersSoADevice(const SiStripClustersSoADevice &&) = delete;
SiStripClustersSoADevice &operator=(const SiStripClustersSoADevice &&) = delete;
SiStripClustersSoADevice(SiStripClustersSoADevice &&) = default;
SiStripClustersSoADevice &operator=(SiStripClustersSoADevice &&) = default;

uint32_t maxClusters() const { return maxClusters_; }

private:
uint32_t maxClusters_;
};

#endif
28 changes: 28 additions & 0 deletions CUDADataFormats/SiStripCluster/interface/SiStripClustersSoAHost.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,28 @@
#ifndef CUDADataFormats_SiStripCluster_interface_SiStripClustersSoAHost_h
#define CUDADataFormats_SiStripCluster_interface_SiStripClustersSoAHost_h

#include "CUDADataFormats/Common/interface/PortableHostCollection.h"
#include "CUDADataFormats/SiStripCluster/interface/SiStripClustersSoAUtilities.h"

class SiStripClustersSoAHost : public cms::cuda::PortableHostCollection<SiStripClustersLayout> {
public:
using cms::cuda::PortableHostCollection<SiStripClustersLayout>::view;
using cms::cuda::PortableHostCollection<SiStripClustersLayout>::const_view;
using cms::cuda::PortableHostCollection<SiStripClustersLayout>::buffer;
using cms::cuda::PortableHostCollection<SiStripClustersLayout>::bufferSize;

SiStripClustersSoAHost() = default;
~SiStripClustersSoAHost() = default;

explicit SiStripClustersSoAHost(uint32_t maxClusters, cudaStream_t stream)
: PortableHostCollection<SiStripClustersLayout>(maxClusters, stream){};

SiStripClustersSoAHost(const SiStripClustersSoAHost &&) = delete;
SiStripClustersSoAHost &operator=(const SiStripClustersSoAHost &&) = delete;
SiStripClustersSoAHost(SiStripClustersSoAHost &&) = default;
SiStripClustersSoAHost &operator=(SiStripClustersSoAHost &&) = default;

private:
};

#endif
27 changes: 27 additions & 0 deletions CUDADataFormats/SiStripCluster/interface/SiStripClustersSoAUtilities.h
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -0,0 +1,27 @@
#ifndef CUDADataFormats_SiStripCluster_interface_SiStripClustersSoAUtilities_h
#define CUDADataFormats_SiStripCluster_interface_SiStripClustersSoAUtilities_h

#include "DataFormats/SiStripCluster/interface/SiStripTypes.h"
#include "DataFormats/SoATemplate/interface/SoALayout.h"

struct SiStripClustersSoA {
const static auto maxStripsPerCluster = 768;
using clusterADCsColumn = std::array<uint8_t, maxStripsPerCluster /*768*/>;
GENERATE_SOA_LAYOUT(SiStripClustersSoALayout,
SOA_COLUMN(uint32_t, clusterIndex),
SOA_COLUMN(uint32_t, clusterSize),
SOA_COLUMN(clusterADCsColumn, clusterADCs),
SOA_COLUMN(stripgpu::detId_t, clusterDetId),
SOA_COLUMN(stripgpu::stripId_t, firstStrip),
SOA_COLUMN(bool, trueCluster),
SOA_COLUMN(float, barycenter),
SOA_COLUMN(float, charge),
SOA_SCALAR(uint32_t, nClusters),
SOA_SCALAR(uint32_t, maxClusterSize));
};

using SiStripClustersLayout = typename SiStripClustersSoA::SiStripClustersSoALayout<>;
using SiStripClustersView = typename SiStripClustersSoA::SiStripClustersSoALayout<>::View;
using SiStripClustersConstView = typename SiStripClustersSoA::SiStripClustersSoALayout<>::ConstView;

#endif
59 changes: 0 additions & 59 deletions CUDADataFormats/SiStripCluster/src/SiStripClustersCUDA.cc
View file
Open in desktop

This file was deleted.

3 changes: 2 additions & 1 deletion CUDADataFormats/SiStripCluster/src/classes.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2,7 +2,8 @@
#define CUDADataFormats_SiStripCluster_classes_h

#include "CUDADataFormats/Common/interface/Product.h"
#include "CUDADataFormats/SiStripCluster/interface/SiStripClustersCUDA.h"
#include "CUDADataFormats/SiStripCluster/interface/SiStripClustersSoADevice.h"
#include "CUDADataFormats/SiStripCluster/interface/SiStripClustersSoAHost.h"
#include "DataFormats/Common/interface/Wrapper.h"

#endif
15 changes: 11 additions & 4 deletions CUDADataFormats/SiStripCluster/src/classes_def.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,13 @@
<lcgdict>
<class name="cms::cuda::Product<SiStripClustersCUDADevice>" persistent="false"/>
<class name="edm::Wrapper<cms::cuda::Product<SiStripClustersCUDADevice>>" persistent="false"/>
<class name="SiStripClustersCUDAHost" persistent="false"/>
<class name="edm::Wrapper<SiStripClustersCUDAHost>" persistent="false"/>
<class name="SiStripClustersSoADevice" persistent="false"/>
<class name="edm::Wrapper<SiStripClustersSoADevice>" persistent="false"/>

<class name="cms::cuda::Product<SiStripClustersSoADevice>" persistent="false"/>
<class name="edm::Wrapper<cms::cuda::Product<SiStripClustersSoADevice>>" persistent="false"/>

<class name="SiStripClustersSoAHost" persistent="false"/>
<class name="edm::Wrapper<SiStripClustersSoAHost>" persistent="false"/>

<class name="cms::cuda::Product<SiStripClustersSoAHost>" persistent="false"/>
<class name="edm::Wrapper<cms::cuda::Product<SiStripClustersSoAHost>>" persistent="false"/>
</lcgdict>
7 changes: 3 additions & 4 deletions RecoLocalTracker/SiStripClusterizer/plugins/ClustersFromRawProducerGPU.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -68,7 +68,7 @@ class SiStripClusterizerFromRawGPU final : public edm::stream::EDProducer<edm::E
raw_(sistrip::FED_ID_MAX),
gpuAlgo_(conf.getParameter<edm::ParameterSet>("Clusterizer")),
inputToken_(consumes(conf.getParameter<edm::InputTag>("ProductLabel"))),
outputToken_(produces<cms::cuda::Product<SiStripClustersCUDADevice>>()),
outputToken_(produces<cms::cuda::Product<SiStripClustersSoADevice>>()),
conditionsToken_(esConsumes(edm::ESInputTag{"", conf.getParameter<std::string>("ConditionsLabel")})),
cpuConditionsToken_(esConsumes(edm::ESInputTag{"", conf.getParameter<std::string>("ConditionsLabel")})) {}

Expand All @@ -91,7 +91,6 @@ class SiStripClusterizerFromRawGPU final : public edm::stream::EDProducer<edm::E
// Queues asynchronous data transfers and kernels to the CUDA stream
// returned by cms::cuda::ScopedContextAcquire::stream()
gpuAlgo_.makeAsync(raw_, buffers_, conditions, ctx.stream());

// Destructor of ctx queues a callback to the CUDA stream notifying
// waitingTaskHolder when the queued asynchronous work has finished
}
Expand All @@ -104,7 +103,7 @@ class SiStripClusterizerFromRawGPU final : public edm::stream::EDProducer<edm::E
// OutputData to cms::cuda::Product<OutputData>. cms::cuda::Product<T> stores also
// the current device and the CUDA stream since those will be needed
// in the consumer side.
ctx.emplace(ev, outputToken_, gpuAlgo_.getResults(ctx.stream()));
ctx.emplace(ev, outputToken_, std::move(gpuAlgo_.getResults(ctx.stream())));

for (auto& buf : buffers_)
buf.reset(nullptr);
Expand All @@ -122,7 +121,7 @@ class SiStripClusterizerFromRawGPU final : public edm::stream::EDProducer<edm::E
stripgpu::SiStripRawToClusterGPUKernel gpuAlgo_;

edm::EDGetTokenT<FEDRawDataCollection> inputToken_;
edm::EDPutTokenT<cms::cuda::Product<SiStripClustersCUDADevice>> outputToken_;
edm::EDPutTokenT<cms::cuda::Product<SiStripClustersSoADevice>> outputToken_;
edm::ESGetToken<stripgpu::SiStripClusterizerConditionsGPU, SiStripClusterizerConditionsRcd> conditionsToken_;
edm::ESGetToken<SiStripClusterizerConditions, SiStripClusterizerConditionsRcd> cpuConditionsToken_;
};
Expand Down
23 changes: 13 additions & 10 deletions RecoLocalTracker/SiStripClusterizer/plugins/SiStripClustersFromSOA.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,14 +13,17 @@
#include "FWCore/Framework/interface/ESHandle.h"
#include "FWCore/MessageLogger/interface/MessageLogger.h"

#include "CUDADataFormats/SiStripCluster/interface/SiStripClustersCUDA.h"
#include "CUDADataFormats/SiStripCluster/interface/SiStripClustersSoAHost.h"

#include "CUDADataFormats/Common/interface/Product.h"
#include "FWCore/Framework/interface/ConsumesCollector.h"
#include <iostream>
#include <memory>

class SiStripClustersFromSOA final : public edm::stream::EDProducer<> {
public:
explicit SiStripClustersFromSOA(const edm::ParameterSet& conf)
: inputToken_(consumes<SiStripClustersCUDAHost>(conf.getParameter<edm::InputTag>("ProductLabel"))),
: inputToken_(consumes(conf.getParameter<edm::InputTag>("ProductLabel"))),
outputToken_(produces<edmNew::DetSetVector<SiStripCluster>>()) {}

static void fillDescriptions(edm::ConfigurationDescriptions& descriptions) {
Expand All @@ -34,12 +37,12 @@ class SiStripClustersFromSOA final : public edm::stream::EDProducer<> {
void produce(edm::Event& ev, const edm::EventSetup& es) override {
const auto& clust_data = ev.get(inputToken_);

const int nSeedStripsNC = clust_data.nClusters();
const auto clusterSize = clust_data.clusterSize().get();
const auto clusterADCs = clust_data.clusterADCs().get();
const auto detIDs = clust_data.clusterDetId().get();
const auto stripIDs = clust_data.firstStrip().get();
const auto trueCluster = clust_data.trueCluster().get();
const int nSeedStripsNC = clust_data->nClusters();
const auto clusterSize = clust_data->clusterSize();
const auto clusterADCs = clust_data->clusterADCs();
const auto detIDs = clust_data->clusterDetId();
const auto stripIDs = clust_data->firstStrip();
const auto trueCluster = clust_data->trueCluster();

const unsigned int initSeedStripsSize = 15000;

Expand All @@ -62,7 +65,7 @@ class SiStripClustersFromSOA final : public edm::stream::EDProducer<> {
adcs.reserve(size);

for (uint32_t j = 0; j < size; ++j) {
adcs.push_back(clusterADCs[i + j * nSeedStripsNC]);
adcs.push_back(clusterADCs[i][j]);
}
record.push_back(SiStripCluster(firstStrip, std::move(adcs)));
}
Expand All @@ -75,7 +78,7 @@ class SiStripClustersFromSOA final : public edm::stream::EDProducer<> {
}

private:
edm::EDGetTokenT<SiStripClustersCUDAHost> inputToken_;
edm::EDGetTokenT<SiStripClustersSoAHost> inputToken_;
edm::EDPutTokenT<edmNew::DetSetVector<SiStripCluster>> outputToken_;
};

Expand Down
Loading