Skip to content

Commit

Permalink
Trackertraits Made Portable (cms-sw#11)
Browse files Browse the repository at this point in the history 
* Portable data formats with Tracker Traits

Co-authored-by: Breno Orzari <breno.orzari@hotmail.com>
Co-authored-by: Dimitris Papagiannis <nothingface0@gmail.com>
Co-authored-by: Dimitris Papagiannis <d.papag.a@gmail.com>
Co-authored-by: borzari <40498845+borzari@users.noreply.github.com>
  • Loading branch information
@AdrianoDee @borzari @nothingface0
5 people authored Dec 5, 2022
1 parent ef3151a commit 73b8e2c
Show file tree
Hide file tree
Showing 100 changed files with 3,890 additions and 2,640 deletions.
55 changes: 17 additions & 38 deletions CUDADataFormats/SiPixelCluster/interface/SiPixelClustersCUDA.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -5,16 +5,29 @@
#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<>;
using SiPixelClustersCUDASOAView = SiPixelClustersCUDALayout<>::View;
using SiPixelClustersCUDASOAConstView = SiPixelClustersCUDALayout<>::ConstView;

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

SiPixelClustersCUDA(const SiPixelClustersCUDA &) = delete;
SiPixelClustersCUDA &operator=(const SiPixelClustersCUDA &) = delete;
explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
: PortableDeviceCollection<SiPixelClustersCUDALayout<>>(maxModules + 1, stream) {}

SiPixelClustersCUDA(SiPixelClustersCUDA &&) = default;
SiPixelClustersCUDA &operator=(SiPixelClustersCUDA &&) = default;

Expand All @@ -26,41 +39,7 @@ 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(); }

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

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);
// }
45 changes: 34 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,33 @@
#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 +44,23 @@ 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 73b8e2c

Please sign in to comment.