Skip to content

Commit

Permalink
Making pixel tracks consuming new rec hits structures
Browse files Browse the repository at this point in the history
  • Loading branch information
AdrianoDee committed Nov 21, 2022
1 parent 617a49f commit 65eb849
Show file tree
Hide file tree
Showing 42 changed files with 761 additions and 627 deletions.
91 changes: 55 additions & 36 deletions CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -5,16 +5,35 @@
#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"

#include "DataFormats/SoATemplate/interface/SoALayout.h"
#include "CUDADataFormats/Common/interface/PortableDeviceCollection.h"

#include <cuda_runtime.h>

class SiPixelClustersCUDA {
GENERATE_SOA_LAYOUT(SiPixelClustersCUDALayout,
SOA_COLUMN(uint32_t, moduleStart),
SOA_COLUMN(uint32_t, clusInModule),
SOA_COLUMN(uint32_t, moduleId),
SOA_COLUMN(uint32_t, clusModuleStart))

using SiPixelClustersCUDASoA = SiPixelClustersCUDALayout<>;

class SiPixelClustersCUDA : public cms::cuda::PortableDeviceCollection<SiPixelClustersCUDALayout<>> {
public:
SiPixelClustersCUDA() = default;
explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream);
// explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream);
explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
: PortableDeviceCollection<SiPixelClustersCUDALayout<>>(maxModules + 1, stream) {}
~SiPixelClustersCUDA() = default;

SiPixelClustersCUDA(const SiPixelClustersCUDA &) = delete;
SiPixelClustersCUDA &operator=(const SiPixelClustersCUDA &) = delete;
// // Restrict view
// using RestrictConstView =
// Layout::ConstViewTemplate<cms::soa::RestrictQualify::enabled, cms::soa::RangeChecking::disabled>;
//
// RestrictConstView restrictConstView() const { return RestrictConstView(layout()); }

// SiPixelClustersCUDA(const SiPixelClustersCUDA &) = delete;
// SiPixelClustersCUDA &operator=(const SiPixelClustersCUDA &) = delete;
SiPixelClustersCUDA(SiPixelClustersCUDA &&) = default;
SiPixelClustersCUDA &operator=(SiPixelClustersCUDA &&) = default;

Expand All @@ -26,40 +45,40 @@ class SiPixelClustersCUDA {
uint32_t nClusters() const { return nClusters_h; }
int32_t offsetBPIX2() const { return offsetBPIX2_h; }

uint32_t *moduleStart() { return moduleStart_d.get(); }
uint32_t *clusInModule() { return clusInModule_d.get(); }
uint32_t *moduleId() { return moduleId_d.get(); }
uint32_t *clusModuleStart() { return clusModuleStart_d.get(); }

uint32_t const *moduleStart() const { return moduleStart_d.get(); }
uint32_t const *clusInModule() const { return clusInModule_d.get(); }
uint32_t const *moduleId() const { return moduleId_d.get(); }
uint32_t const *clusModuleStart() const { return clusModuleStart_d.get(); }

class SiPixelClustersCUDASOAView {
public:
__device__ __forceinline__ uint32_t moduleStart(int i) const { return __ldg(moduleStart_ + i); }
__device__ __forceinline__ uint32_t clusInModule(int i) const { return __ldg(clusInModule_ + i); }
__device__ __forceinline__ uint32_t moduleId(int i) const { return __ldg(moduleId_ + i); }
__device__ __forceinline__ uint32_t clusModuleStart(int i) const { return __ldg(clusModuleStart_ + i); }

uint32_t const *moduleStart_;
uint32_t const *clusInModule_;
uint32_t const *moduleId_;
uint32_t const *clusModuleStart_;
};

SiPixelClustersCUDASOAView const *view() const { return view_d.get(); }
// uint32_t *moduleStart() { return moduleStart_d.get(); }
// uint32_t *clusInModule() { return clusInModule_d.get(); }
// uint32_t *moduleId() { return moduleId_d.get(); }
// uint32_t *clusModuleStart() { return clusModuleStart_d.get(); }
//
// uint32_t const *moduleStart() const { return moduleStart_d.get(); }
// uint32_t const *clusInModule() const { return clusInModule_d.get(); }
// uint32_t const *moduleId() const { return moduleId_d.get(); }
// uint32_t const *clusModuleStart() const { return clusModuleStart_d.get(); }
//
// class SiPixelClustersCUDASOAView {
// public:
// __device__ __forceinline__ uint32_t moduleStart(int i) const { return __ldg(moduleStart_ + i); }
// __device__ __forceinline__ uint32_t clusInModule(int i) const { return __ldg(clusInModule_ + i); }
// __device__ __forceinline__ uint32_t moduleId(int i) const { return __ldg(moduleId_ + i); }
// __device__ __forceinline__ uint32_t clusModuleStart(int i) const { return __ldg(clusModuleStart_ + i); }
//
// uint32_t const *moduleStart_;
// uint32_t const *clusInModule_;
// uint32_t const *moduleId_;
// uint32_t const *clusModuleStart_;
// };
//
// SiPixelClustersCUDASOAView const *view() const { return view_d.get(); }

private:
cms::cuda::device::unique_ptr<uint32_t[]> moduleStart_d; // index of the first pixel of each module
cms::cuda::device::unique_ptr<uint32_t[]> clusInModule_d; // number of clusters found in each module
cms::cuda::device::unique_ptr<uint32_t[]> moduleId_d; // module id of each module

// originally from rechits
cms::cuda::device::unique_ptr<uint32_t[]> clusModuleStart_d; // index of the first cluster of each module

cms::cuda::device::unique_ptr<SiPixelClustersCUDASOAView> view_d; // "me" pointer
// cms::cuda::device::unique_ptr<uint32_t[]> moduleStart_d; // index of the first pixel of each module
// cms::cuda::device::unique_ptr<uint32_t[]> clusInModule_d; // number of clusters found in each module
// cms::cuda::device::unique_ptr<uint32_t[]> moduleId_d; // module id of each module
//
// // originally from rechits
// cms::cuda::device::unique_ptr<uint32_t[]> clusModuleStart_d; // index of the first cluster of each module
//
// cms::cuda::device::unique_ptr<SiPixelClustersCUDASOAView> view_d; // "me" pointer

uint32_t nClusters_h = 0;
int32_t offsetBPIX2_h = 0;
Expand Down
38 changes: 19 additions & 19 deletions CUDADataFormats/SiPixelCluster/src/SiPixelClustersCUDA.cc
View file
Original file line number Diff line number Diff line change
@@ -1,19 +1,19 @@
#include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"

SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
: moduleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)),
clusInModule_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules, stream)),
moduleId_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules, stream)),
clusModuleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)) {
auto view = cms::cuda::make_host_unique<SiPixelClustersCUDASOAView>(stream);
view->moduleStart_ = moduleStart_d.get();
view->clusInModule_ = clusInModule_d.get();
view->moduleId_ = moduleId_d.get();
view->clusModuleStart_ = clusModuleStart_d.get();

view_d = cms::cuda::make_device_unique<SiPixelClustersCUDASOAView>(stream);
cms::cuda::copyAsync(view_d, view, stream);
}
// #include "CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h"
// #include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h"
// #include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
// #include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
//
// SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
// : moduleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)),
// clusInModule_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules, stream)),
// moduleId_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules, stream)),
// clusModuleStart_d(cms::cuda::make_device_unique<uint32_t[]>(maxModules + 1, stream)) {
// auto view = cms::cuda::make_host_unique<SiPixelClustersCUDASOAView>(stream);
// view->moduleStart_ = moduleStart_d.get();
// view->clusInModule_ = clusInModule_d.get();
// view->moduleId_ = moduleId_d.get();
// view->clusModuleStart_ = clusModuleStart_d.get();
//
// view_d = cms::cuda::make_device_unique<SiPixelClustersCUDASOAView>(stream);
// cms::cuda::copyAsync(view_d, view, stream);
// }
47 changes: 36 additions & 11 deletions CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDA.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,17 +6,34 @@
#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h"
#include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDASOAView.h"
// #include "CUDADataFormats/SiPixelDigi/interface/SiPixelDigisCUDASOAView.h"
#include "CUDADataFormats/Common/interface/PortableDeviceCollection.h"
#include "DataFormats/SoATemplate/interface/SoALayout.h"

class SiPixelDigisCUDA {

GENERATE_SOA_LAYOUT(SiPixelDigisSoALayout,
SOA_COLUMN(int32_t, clus),
SOA_COLUMN(uint32_t, pdigi),
SOA_COLUMN(uint32_t, rawIdArr),
SOA_COLUMN(uint16_t, adc),
SOA_COLUMN(uint16_t, xx),
SOA_COLUMN(uint16_t, yy),
SOA_COLUMN(uint16_t, moduleId))

using SiPixelDigisCUDASOA = SiPixelDigisSoALayout<>;
using SiPixelDigisCUDASOAView = SiPixelDigisCUDASOA::View;
using SiPixelDigisCUDASOAConstView = SiPixelDigisCUDASOA::ConstView;

class SiPixelDigisCUDA : public cms::cuda::PortableDeviceCollection<SiPixelDigisSoALayout<>> {
public:
using StoreType = uint16_t;
// using StoreType = uint16_t;
SiPixelDigisCUDA() = default;
explicit SiPixelDigisCUDA(size_t maxFedWords, cudaStream_t stream);
explicit SiPixelDigisCUDA(size_t maxFedWords, cudaStream_t stream)
: PortableDeviceCollection<SiPixelDigisSoALayout<>>(maxFedWords + 1, stream) {}
~SiPixelDigisCUDA() = default;

SiPixelDigisCUDA(const SiPixelDigisCUDA &) = delete;
SiPixelDigisCUDA &operator=(const SiPixelDigisCUDA &) = delete;
// SiPixelDigisCUDA(const SiPixelDigisCUDA &) = delete;
// SiPixelDigisCUDA &operator=(const SiPixelDigisCUDA &) = delete;
SiPixelDigisCUDA(SiPixelDigisCUDA &&) = default;
SiPixelDigisCUDA &operator=(SiPixelDigisCUDA &&) = default;

Expand All @@ -28,16 +45,24 @@ class SiPixelDigisCUDA {
uint32_t nModules() const { return nModules_h; }
uint32_t nDigis() const { return nDigis_h; }

cms::cuda::host::unique_ptr<StoreType[]> copyAllToHostAsync(cudaStream_t stream) const;
// cms::cuda::host::unique_ptr<StoreType[]> copyAllToHostAsync(cudaStream_t stream) const;

cms::cuda::host::unique_ptr<std::byte[]> copyAllToHostAsync(cudaStream_t stream) const {
// Copy to a host buffer the host-device shared part (m_hostDeviceLayout).
auto ret = cms::cuda::make_host_unique<std::byte[]>(bufferSize(), stream);
cudaCheck(cudaMemcpyAsync(
ret.get(), buffer().get(), bufferSize(), cudaMemcpyDeviceToHost, stream));
return ret;
}

SiPixelDigisCUDASOAView view() { return m_view; }
SiPixelDigisCUDASOAView const view() const { return m_view; }
// SiPixelDigisCUDASOAView view() { return m_view; }
// SiPixelDigisCUDASOAView const view() const { return m_view; }

private:
// These are consumed by downstream device code
cms::cuda::device::unique_ptr<StoreType[]> m_store;
// cms::cuda::device::unique_ptr<StoreType[]> m_store;

SiPixelDigisCUDASOAView m_view;
// SiPixelDigisCUDASOAView m_view;

uint32_t nModules_h = 0;
uint32_t nDigis_h = 0;
Expand Down
Loading

0 comments on commit 65eb849

Please sign in to comment.