Skip to content

Commit

Permalink
[cudadev] Ran clang-format to format code.
Browse files Browse the repository at this point in the history
  • Loading branch information
@ericcano
ericcano committed Jan 17, 2022
1 parent f72044e commit a839a44
Show file tree
Hide file tree
Showing 17 changed files with 919 additions and 922 deletions.
5 changes: 2 additions & 3 deletions src/cudadev/CUDADataFormats/SiPixelClustersCUDA.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -3,10 +3,9 @@
#include "CUDACore/host_unique_ptr.h"
#include "CUDADataFormats/SiPixelClustersCUDA.h"

SiPixelClustersCUDA::SiPixelClustersCUDA(): data_d(), deviceLayout_(data_d.get(), 0), deviceView_(deviceLayout_) {}
SiPixelClustersCUDA::SiPixelClustersCUDA() : data_d(), deviceLayout_(data_d.get(), 0), deviceView_(deviceLayout_) {}

SiPixelClustersCUDA::SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream)
: data_d(cms::cuda::make_device_unique<std::byte[]>(DeviceLayout::computeDataSize(maxModules), stream)),
deviceLayout_(data_d.get(), maxModules),
deviceView_(deviceLayout_)
{}
deviceView_(deviceLayout_) {}
67 changes: 31 additions & 36 deletions src/cudadev/CUDADataFormats/SiPixelClustersCUDA.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -12,45 +12,40 @@
class SiPixelClustersCUDA {
public:
GENERATE_SOA_LAYOUT(DeviceLayoutTemplate,
SOA_COLUMN(uint32_t, moduleStart), // index of the first pixel of each module
SOA_COLUMN(uint32_t, clusInModule), // number of clusters found in each module
SOA_COLUMN(uint32_t, moduleId), // module id of each module

// originally from rechits
SOA_COLUMN(uint32_t, clusModuleStart) // index of the first cluster of each module
)

SOA_COLUMN(uint32_t, moduleStart), // index of the first pixel of each module
SOA_COLUMN(uint32_t, clusInModule), // number of clusters found in each module
SOA_COLUMN(uint32_t, moduleId), // module id of each module

// originally from rechits
SOA_COLUMN(uint32_t, clusModuleStart)) // index of the first cluster of each module

// We use all defaults for the template parameters.
using DeviceLayout = DeviceLayoutTemplate<>;

GENERATE_SOA_VIEW(DeviceViewTemplate,
SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceLayout, deviceLayout)),
SOA_VIEW_VALUE_LIST(
SOA_VIEW_VALUE(deviceLayout, moduleStart), // index of the first pixel of each module
SOA_VIEW_VALUE(deviceLayout, clusInModule), // number of clusters found in each module
SOA_VIEW_VALUE(deviceLayout, moduleId), // module id of each module

// originally from rechits
SOA_VIEW_VALUE(deviceLayout, clusModuleStart) // index of the first cluster of each module
)
)

GENERATE_SOA_VIEW(
DeviceViewTemplate,
SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceLayout, deviceLayout)),
SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(deviceLayout, moduleStart), // index of the first pixel of each module
SOA_VIEW_VALUE(deviceLayout, clusInModule), // number of clusters found in each module
SOA_VIEW_VALUE(deviceLayout, moduleId), // module id of each module

// originally from rechits
SOA_VIEW_VALUE(deviceLayout, clusModuleStart))) // index of the first cluster of each module

using DeviceView = DeviceViewTemplate<>;

GENERATE_SOA_CONST_VIEW(DeviceConstViewTemplate,
SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceView, deviceView)),
SOA_VIEW_VALUE_LIST(
SOA_VIEW_VALUE(deviceView, moduleStart), // index of the first pixel of each module
SOA_VIEW_VALUE(deviceView, clusInModule), // number of clusters found in each module
SOA_VIEW_VALUE(deviceView, moduleId), // module id of each module

// originally from rechits
SOA_VIEW_VALUE(deviceView, clusModuleStart) // index of the first cluster of each module
)
)


GENERATE_SOA_CONST_VIEW(
DeviceConstViewTemplate,
SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceView, deviceView)),
SOA_VIEW_VALUE_LIST(SOA_VIEW_VALUE(deviceView, moduleStart), // index of the first pixel of each module
SOA_VIEW_VALUE(deviceView, clusInModule), // number of clusters found in each module
SOA_VIEW_VALUE(deviceView, moduleId), // module id of each module

// originally from rechits
SOA_VIEW_VALUE(deviceView, clusModuleStart))) // index of the first cluster of each module

using DeviceConstView = DeviceConstViewTemplate<>;

explicit SiPixelClustersCUDA();
explicit SiPixelClustersCUDA(size_t maxModules, cudaStream_t stream);
~SiPixelClustersCUDA() = default;
Expand All @@ -77,10 +72,10 @@ class SiPixelClustersCUDA {
DeviceConstView view() const { return DeviceConstView(deviceView_); }

private:
cms::cuda::device::unique_ptr<std::byte[]> data_d; // Single SoA storage
cms::cuda::device::unique_ptr<std::byte[]> data_d; // Single SoA storage
DeviceLayout deviceLayout_;
DeviceView deviceView_;

uint32_t nClusters_h = 0;
};

Expand Down
44 changes: 23 additions & 21 deletions src/cudadev/CUDADataFormats/SiPixelDigisCUDA.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -6,28 +6,22 @@

SiPixelDigisCUDA::SiPixelDigisCUDA(size_t maxFedWords, cudaStream_t stream)
: data_d(cms::cuda::make_device_unique<std::byte[]>(
DeviceOnlyLayout::computeDataSize(maxFedWords) +
HostDeviceLayout::computeDataSize(maxFedWords),
stream)),
DeviceOnlyLayout::computeDataSize(maxFedWords) + HostDeviceLayout::computeDataSize(maxFedWords), stream)),
deviceOnlyLayout_d(data_d.get(), maxFedWords),
hostDeviceLayout_d(deviceOnlyLayout_d.soaMetadata().nextByte(), maxFedWords),
deviceFullView_(deviceOnlyLayout_d, hostDeviceLayout_d),
devicePixelConstView_(deviceFullView_)
{}
devicePixelConstView_(deviceFullView_) {}

SiPixelDigisCUDA::SiPixelDigisCUDA()
: data_d(),deviceOnlyLayout_d(), hostDeviceLayout_d(), deviceFullView_(), devicePixelConstView_()
{}
: data_d(), deviceOnlyLayout_d(), hostDeviceLayout_d(), deviceFullView_(), devicePixelConstView_() {}

SiPixelDigisCUDA::HostStore::HostStore()
: data_h(), hostLayout_(nullptr, 0), hostView_(hostLayout_)
{}
SiPixelDigisCUDA::HostStore::HostStore() : data_h(), hostLayout_(nullptr, 0), hostView_(hostLayout_) {}

SiPixelDigisCUDA::HostStore::HostStore(size_t maxFedWords, cudaStream_t stream)
: data_h(cms::cuda::make_host_unique<std::byte[]>(SiPixelDigisCUDA::HostDeviceLayout::computeDataSize(maxFedWords), stream)),
hostLayout_(data_h.get(), maxFedWords),
hostView_(hostLayout_)
{}
: data_h(cms::cuda::make_host_unique<std::byte[]>(SiPixelDigisCUDA::HostDeviceLayout::computeDataSize(maxFedWords),
stream)),
hostLayout_(data_h.get(), maxFedWords),
hostView_(hostLayout_) {}

void SiPixelDigisCUDA::HostStore::reset() {
hostLayout_ = HostDeviceLayout();
Expand All @@ -38,7 +32,8 @@ void SiPixelDigisCUDA::HostStore::reset() {
cms::cuda::host::unique_ptr<uint16_t[]> SiPixelDigisCUDA::adcToHostAsync(cudaStream_t stream) const {
auto ret = cms::cuda::make_host_unique<uint16_t[]>(nDigis(), stream);
// TODO: this is downgraded from cms::cuda::copyAsync as we copy data from within a block but not the full block.
cudaCheck(cudaMemcpyAsync(ret.get(), deviceFullView_.adc(), nDigis() * sizeof(decltype(ret[0])), cudaMemcpyDeviceToHost, stream));
cudaCheck(cudaMemcpyAsync(
ret.get(), deviceFullView_.adc(), nDigis() * sizeof(decltype(ret[0])), cudaMemcpyDeviceToHost, stream));
return ret;
}

Expand All @@ -49,13 +44,20 @@ SiPixelDigisCUDA::HostStore SiPixelDigisCUDA::dataToHostAsync(cudaStream_t strea
HostStore ret(nDigis(), stream);
auto rhlsm = ret.hostLayout_.soaMetadata();
auto hdlsm_d = hostDeviceLayout_d.soaMetadata();
cudaCheck(cudaMemcpyAsync(rhlsm.addressOf_adc(), hdlsm_d.addressOf_adc(), nDigis_h * sizeof(*rhlsm.addressOf_adc()),
cudaMemcpyDeviceToHost, stream));
cudaCheck(cudaMemcpyAsync(rhlsm.addressOf_adc(),
hdlsm_d.addressOf_adc(),
nDigis_h * sizeof(*rhlsm.addressOf_adc()),
cudaMemcpyDeviceToHost,
stream));
// Copy the other columns, realigning the data in shorter arrays. clus is the first but all 3 columns (clus, pdigis, rawIdArr) have
// the same geometry.
cudaCheck(cudaMemcpy2DAsync(rhlsm.addressOf_clus(), rhlsm.clusPitch(),
hdlsm_d.addressOf_clus(), hdlsm_d.clusPitch(),
3 /* rows */,
nDigis() * sizeof(decltype (*ret.hostView_.clus())), cudaMemcpyDeviceToHost, stream));
cudaCheck(cudaMemcpy2DAsync(rhlsm.addressOf_clus(),
rhlsm.clusPitch(),
hdlsm_d.addressOf_clus(),
hdlsm_d.clusPitch(),
3 /* rows */,
nDigis() * sizeof(decltype(*ret.hostView_.clus())),
cudaMemcpyDeviceToHost,
stream));
return ret;
}
127 changes: 62 additions & 65 deletions src/cudadev/CUDADataFormats/SiPixelDigisCUDA.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,75 +11,71 @@

class SiPixelDigisCUDA {
public:
GENERATE_SOA_LAYOUT(DeviceOnlyLayoutTemplate,
/* These are consumed by downstream device code */
SOA_COLUMN(uint16_t, xx), /* local coordinates of each pixel */
SOA_COLUMN(uint16_t, yy), /* */
SOA_COLUMN(uint16_t, moduleInd) /* module id of each pixel */
GENERATE_SOA_LAYOUT(
DeviceOnlyLayoutTemplate,
/* These are consumed by downstream device code */
SOA_COLUMN(uint16_t, xx), /* local coordinates of each pixel */
SOA_COLUMN(uint16_t, yy), /* */
SOA_COLUMN(uint16_t, moduleInd) /* module id of each pixel */
)

using DeviceOnlyLayout = DeviceOnlyLayoutTemplate<>;

GENERATE_SOA_LAYOUT(HostDeviceLayoutTemplate,
/* These are also transferred to host (see HostDataView) */
SOA_COLUMN(uint16_t, adc), /* ADC of each pixel */
SOA_COLUMN(int32_t, clus), /* cluster id of each pixel */
/* These are for CPU output; should we (eventually) place them to a */
/* separate product? */
SOA_COLUMN(uint32_t, pdigi), /* packed digi (row, col, adc) of each pixel */
SOA_COLUMN(uint32_t, rawIdArr) /* DetId of each pixel */

GENERATE_SOA_LAYOUT(
HostDeviceLayoutTemplate,
/* These are also transferred to host (see HostDataView) */
SOA_COLUMN(uint16_t, adc), /* ADC of each pixel */
SOA_COLUMN(int32_t, clus), /* cluster id of each pixel */
/* These are for CPU output; should we (eventually) place them to a */
/* separate product? */
SOA_COLUMN(uint32_t, pdigi), /* packed digi (row, col, adc) of each pixel */
SOA_COLUMN(uint32_t, rawIdArr) /* DetId of each pixel */
)

using HostDeviceLayout = HostDeviceLayoutTemplate<>;

GENERATE_SOA_VIEW(HostDeviceViewTemplate,
SOA_VIEW_LAYOUT_LIST(
SOA_VIEW_LAYOUT(HostDeviceLayout, hostDevice)
),
SOA_VIEW_VALUE_LIST(
SOA_VIEW_VALUE(hostDevice, adc), /* ADC of each pixel */
SOA_VIEW_VALUE(hostDevice, clus), /* cluster id of each pixel */
SOA_VIEW_VALUE(hostDevice, pdigi), /* packed digi (row, col, adc) of each pixel */
SOA_VIEW_VALUE(hostDevice, rawIdArr) /* DetId of each pixel */
)
)

SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(HostDeviceLayout, hostDevice)),
SOA_VIEW_VALUE_LIST(
SOA_VIEW_VALUE(hostDevice, adc), /* ADC of each pixel */
SOA_VIEW_VALUE(hostDevice, clus), /* cluster id of each pixel */
SOA_VIEW_VALUE(hostDevice, pdigi), /* packed digi (row, col, adc) of each pixel */
SOA_VIEW_VALUE(hostDevice,
rawIdArr) /* DetId of each pixel */
))

using HostDeviceView = HostDeviceViewTemplate<>;

GENERATE_SOA_VIEW(DeviceFullViewTemplate,
SOA_VIEW_LAYOUT_LIST(
SOA_VIEW_LAYOUT(DeviceOnlyLayout, deviceOnly),
SOA_VIEW_LAYOUT(HostDeviceLayout, hostDevice)
),
SOA_VIEW_VALUE_LIST(
SOA_VIEW_VALUE(deviceOnly, xx), /* local coordinates of each pixel */
SOA_VIEW_VALUE(deviceOnly, yy), /* */
SOA_VIEW_VALUE(deviceOnly, moduleInd),/* module id of each pixel */
SOA_VIEW_VALUE(hostDevice, adc), /* ADC of each pixel */
SOA_VIEW_VALUE(hostDevice, clus), /* cluster id of each pixel */
SOA_VIEW_VALUE(hostDevice, pdigi), /* packed digi (row, col, adc) of each pixel */
SOA_VIEW_VALUE(hostDevice, rawIdArr) /* DetId of each pixel */
)
)


GENERATE_SOA_VIEW(
DeviceFullViewTemplate,
SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceOnlyLayout, deviceOnly),
SOA_VIEW_LAYOUT(HostDeviceLayout, hostDevice)),
SOA_VIEW_VALUE_LIST(
SOA_VIEW_VALUE(deviceOnly, xx), /* local coordinates of each pixel */
SOA_VIEW_VALUE(deviceOnly, yy), /* */
SOA_VIEW_VALUE(deviceOnly, moduleInd), /* module id of each pixel */
SOA_VIEW_VALUE(hostDevice, adc), /* ADC of each pixel */
SOA_VIEW_VALUE(hostDevice, clus), /* cluster id of each pixel */
SOA_VIEW_VALUE(hostDevice, pdigi), /* packed digi (row, col, adc) of each pixel */
SOA_VIEW_VALUE(hostDevice, rawIdArr) /* DetId of each pixel */
))

using DeviceFullView = DeviceFullViewTemplate<>;

/* Device pixel view: this is a second generation view (view from view) */
GENERATE_SOA_CONST_VIEW(DevicePixelConstViewTemplate,
/* We get out data from the DeviceFullView */
SOA_VIEW_LAYOUT_LIST(
SOA_VIEW_LAYOUT(DeviceFullView, deviceFullView)
),
/* These are consumed by downstream device code */
SOA_VIEW_VALUE_LIST(
SOA_VIEW_VALUE(deviceFullView, xx), /* local coordinates of each pixel */
SOA_VIEW_VALUE(deviceFullView, yy), /* */
SOA_VIEW_VALUE(deviceFullView, moduleInd), /* module id of each pixel */
SOA_VIEW_VALUE(deviceFullView, adc), /* ADC of each pixel */
SOA_VIEW_VALUE(deviceFullView, clus) /* cluster id of each pixel */
)
)

GENERATE_SOA_CONST_VIEW(
DevicePixelConstViewTemplate,
/* We get out data from the DeviceFullView */
SOA_VIEW_LAYOUT_LIST(SOA_VIEW_LAYOUT(DeviceFullView, deviceFullView)),
/* These are consumed by downstream device code */
SOA_VIEW_VALUE_LIST(
SOA_VIEW_VALUE(deviceFullView, xx), /* local coordinates of each pixel */
SOA_VIEW_VALUE(deviceFullView, yy), /* */
SOA_VIEW_VALUE(deviceFullView, moduleInd), /* module id of each pixel */
SOA_VIEW_VALUE(deviceFullView, adc), /* ADC of each pixel */
SOA_VIEW_VALUE(deviceFullView, clus) /* cluster id of each pixel */
))

using DevicePixelConstView = DevicePixelConstViewTemplate<>;

explicit SiPixelDigisCUDA();
Expand Down Expand Up @@ -117,27 +113,28 @@ class SiPixelDigisCUDA {

class HostStore {
friend SiPixelDigisCUDA;

public:
HostStore();
const SiPixelDigisCUDA::HostDeviceView view() { return hostView_; }
void reset();

private:
HostStore(size_t maxFedWords, cudaStream_t stream);
cms::cuda::host::unique_ptr<std::byte[]> data_h;
HostDeviceLayout hostLayout_;
HostDeviceView hostView_;

};
HostStore dataToHostAsync(cudaStream_t stream) const;

// Special copy for validation
cms::cuda::host::unique_ptr<uint16_t[]> adcToHostAsync(cudaStream_t stream) const;
// Special copy for validation
cms::cuda::host::unique_ptr<uint16_t[]> adcToHostAsync(cudaStream_t stream) const;

const DevicePixelConstView& pixelConstView() const { return devicePixelConstView_; }
const DevicePixelConstView &pixelConstView() const { return devicePixelConstView_; }

private:
// These are consumed by downstream device code
cms::cuda::device::unique_ptr<std::byte[]> data_d; // Single SoA storage
cms::cuda::device::unique_ptr<std::byte[]> data_d; // Single SoA storage
DeviceOnlyLayout deviceOnlyLayout_d;
HostDeviceLayout hostDeviceLayout_d;
DeviceFullView deviceFullView_;
Expand Down
Loading

0 comments on commit a839a44

Please sign in to comment.