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

Multigpu hcal #498

Closed
wants to merge 6 commits into from
Closed

Multigpu hcal #498

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
9ca193e
make scratch use caching alloc
vkhristenko Jun 19, 2020
c7f9392
use caching allocator for input cpu
vkhristenko Jun 19, 2020
a84d731
using allocators for EventHFilter/HcalRawToDigi
vkhristenko Jul 2, 2020
d05f147
hcal/mahi/gpu switch to new input and use allocator for scratch
vkhristenko Jul 3, 2020
8ee109b
hcal/mahi/gpu move pulse offsets to use esproducer and all dev memory…
vkhristenko Jul 3, 2020
9e67a22
remove old
vkhristenko Jul 4, 2020
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
16 changes: 16 additions & 0 deletions CUDADataFormats/HcalCommon/interface/Common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,6 +4,7 @@
#include <vector>

#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"

namespace hcal {
namespace common {
Expand All @@ -13,6 +14,7 @@ namespace hcal {

struct Vec {};
struct Ptr {};
struct DevPtr {};

} // namespace tags

Expand All @@ -24,6 +26,11 @@ namespace hcal {
uint32_t size;
};

template<>
struct AddSize<tags::DevPtr> {
uint32_t size;
};

struct ViewStoragePolicy {
using TagType = tags::Ptr;

Expand All @@ -33,6 +40,15 @@ namespace hcal {
};
};

struct DevStoragePolicy {
using TagType = tags::DevPtr;

template<typename T>
struct StorageSelector {
using type = cms::cuda::device::unique_ptr<T[]>;
};
};

template <template <typename> typename Allocator = std::allocator>
struct VecStoragePolicy {
using TagType = tags::Vec;
Expand Down
2 changes: 2 additions & 0 deletions CUDADataFormats/HcalCommon/src/classes_def.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,6 +1,8 @@
<lcgdict>
<enum name="hcal::common::tags::Ptr"/>
<enum name="hcal::common::tags::Vec"/>
<enum name="hcal::common::tags::DevPtr"/>
<class name="hcal::common::AddSize<hcal::common::tags::Ptr>"/>
<class name="hcal::common::AddSize<hcal::common::tags::Vec>"/>
<class name="hcal::common::AddSize<hcal::common::tags::DevPtr>"/>
</lcgdict>
12 changes: 12 additions & 0 deletions CUDADataFormats/HcalDigi/src/classes_def.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -20,11 +20,23 @@
<class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor4, hcal::common::ViewStoragePolicy>>" persistent="false" />
<class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::ViewStoragePolicy>>" persistent="false" />

<class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::DevStoragePolicy>>" persistent="false" />
<class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor2, hcal::common::DevStoragePolicy>>" persistent="false" />
<class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>>" persistent="false" />
<class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor4, hcal::common::DevStoragePolicy>>" persistent="false" />
<class name="cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::DevStoragePolicy>>" persistent="false" />

<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::ViewStoragePolicy>>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor2, hcal::common::ViewStoragePolicy>>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::ViewStoragePolicy>>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor4, hcal::common::ViewStoragePolicy>>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::ViewStoragePolicy>>>" persistent="false" />

<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor01, hcal::common::DevStoragePolicy>>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor2, hcal::common::DevStoragePolicy>>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor3, hcal::common::DevStoragePolicy>>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor4, hcal::common::DevStoragePolicy>>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::DigiCollection<hcal::Flavor5, hcal::common::DevStoragePolicy>>>" persistent="false" />

<class name="edm::Wrapper<hcal::DigiCollection<hcal::Flavor01, hcal::common::VecStoragePolicy<std::allocator>>>" />
<class name="edm::Wrapper<hcal::DigiCollection<hcal::Flavor2, hcal::common::VecStoragePolicy<std::allocator>>>" />
Expand Down
3 changes: 3 additions & 0 deletions CUDADataFormats/HcalRecHitSoA/src/classes_def.xml
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -7,4 +7,7 @@

<class name="cms::cuda::Product<hcal::RecHitCollection<hcal::common::ViewStoragePolicy>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::RecHitCollection<hcal::common::ViewStoragePolicy>>>" persistent="false" />

<class name="cms::cuda::Product<hcal::RecHitCollection<hcal::common::DevStoragePolicy>>" persistent="false" />
<class name="edm::Wrapper<cms::cuda::Product<hcal::RecHitCollection<hcal::common::DevStoragePolicy>>>" persistent="false" />
</lcgdict>
143 changes: 56 additions & 87 deletions EventFilter/HcalRawToDigi/plugins/DeclsForKernels.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,8 @@
#include "CUDADataFormats/HcalDigi/interface/DigiCollection.h"
#include "HeterogeneousCore/CUDAUtilities/interface/HostAllocator.h"
#include "HeterogeneousCore/CUDAUtilities/interface/cudaCheck.h"
#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h"
#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h"

#include "ElectronicsMappingGPU.h"

Expand Down Expand Up @@ -33,108 +35,75 @@ namespace hcal {
};

struct InputDataCPU {
std::vector<unsigned char, cms::cuda::HostAllocator<unsigned char>> data;
std::vector<uint32_t, cms::cuda::HostAllocator<uint32_t>> offsets;
std::vector<int, cms::cuda::HostAllocator<int>> feds;

void allocate() {
data.resize(utca_nfeds_max * sizeof(unsigned char) * nbytes_per_fed_max);
offsets.resize(utca_nfeds_max, 0);
feds.resize(utca_nfeds_max, 0);
}
cms::cuda::host::unique_ptr<unsigned char[]> data;
cms::cuda::host::unique_ptr<uint32_t[]> offsets;
cms::cuda::host::unique_ptr<int[]> feds;
};

struct OutputDataCPU {
std::vector<uint32_t, cms::cuda::HostAllocator<uint32_t>> nchannels;

void allocate() { nchannels.resize(numOutputCollections); }
cms::cuda::host::unique_ptr<uint32_t[]> nchannels;
};

struct ScratchDataGPU {
// depends on tHE number of output collections
// that is a statically known predefined number!!!
uint32_t *pChannelsCounters = nullptr;

void allocate(ConfigurationParameters const &) {
cudaCheck(cudaMalloc((void **)&pChannelsCounters, sizeof(uint32_t) * numOutputCollections));
}

void deallocate(ConfigurationParameters const &) {
if (pChannelsCounters) {
cudaCheck(cudaFree(pChannelsCounters));
}
}
cms::cuda::device::unique_ptr<uint32_t[]> pChannelsCounters;
};

struct OutputDataGPU {
DigiCollection<Flavor01, common::ViewStoragePolicy> digisF01HE;
DigiCollection<Flavor5, common::ViewStoragePolicy> digisF5HB;
DigiCollection<Flavor3, common::ViewStoragePolicy> digisF3HB;

// qie 11 HE
/*
uint16_t *digisF01HE = nullptr;
uint32_t *idsF01HE = nullptr;

// qie 8 HB
uint16_t *digisF5HB = nullptr;
uint32_t *idsF5HB = nullptr;
uint8_t *npresamplesF5HB = nullptr
*/

void allocate(ConfigurationParameters const &config) {
cudaCheck(
cudaMalloc((void **)&digisF01HE.data,
config.maxChannelsF01HE * sizeof(uint16_t) * compute_stride<Flavor01>(config.nsamplesF01HE)));
cudaCheck(cudaMalloc((void **)&digisF01HE.ids, sizeof(uint32_t) * config.maxChannelsF01HE));

cudaCheck(cudaMalloc((void **)&digisF5HB.data,
config.maxChannelsF5HB * sizeof(uint16_t) * compute_stride<Flavor5>(config.nsamplesF5HB)));
cudaCheck(cudaMalloc((void **)&digisF5HB.ids, sizeof(uint32_t) * config.maxChannelsF5HB));
cudaCheck(cudaMalloc((void **)&digisF5HB.npresamples, sizeof(uint8_t) * config.maxChannelsF5HB));

cudaCheck(cudaMalloc((void **)&digisF3HB.data,
config.maxChannelsF3HB * sizeof(uint16_t) * compute_stride<Flavor3>(config.nsamplesF3HB)));
cudaCheck(cudaMalloc((void **)&digisF3HB.ids, config.maxChannelsF3HB * sizeof(uint32_t)));
}

void deallocate(ConfigurationParameters const &config) {
if (digisF01HE.data) {
cudaCheck(cudaFree(digisF01HE.data));
cudaCheck(cudaFree(digisF01HE.ids));
}

if (digisF5HB.data) {
cudaCheck(cudaFree(digisF5HB.data));
cudaCheck(cudaFree(digisF5HB.ids));
cudaCheck(cudaFree(digisF5HB.npresamples));
}

if (digisF3HB.data) {
cudaCheck(cudaFree(digisF3HB.data));
cudaCheck(cudaFree(digisF3HB.ids));
}
DigiCollection<Flavor01, common::DevStoragePolicy> digisF01HE;
DigiCollection<Flavor5, common::DevStoragePolicy> digisF5HB;
DigiCollection<Flavor3, common::DevStoragePolicy> digisF3HB;

void allocate(ConfigurationParameters const &config, cudaStream_t cudaStream) {
digisF01HE.data = cms::cuda::make_device_unique<uint16_t[]>(
config.maxChannelsF01HE*compute_stride<Flavor01>(config.nsamplesF01HE),
cudaStream
);
//cudaCheck(
// cudaMalloc((void **)&digisF01HE.data,
// config.maxChannelsF01HE * sizeof(uint16_t) * compute_stride<Flavor01>(config.nsamplesF01HE)));
digisF01HE.ids = cms::cuda::make_device_unique<uint32_t[]>(
config.maxChannelsF01HE,
cudaStream
);
//cudaCheck(cudaMalloc((void **)&digisF01HE.ids, sizeof(uint32_t) * config.maxChannelsF01HE));

digisF5HB.data = cms::cuda::make_device_unique<uint16_t[]>(
config.maxChannelsF5HB * compute_stride<Flavor5>(config.nsamplesF5HB),
cudaStream
);
//cudaCheck(cudaMalloc((void **)&digisF5HB.data,
// config.maxChannelsF5HB * sizeof(uint16_t) * compute_stride<Flavor5>(config.nsamplesF5HB)));
digisF5HB.ids = cms::cuda::make_device_unique<uint32_t[]>(
config.maxChannelsF5HB,
cudaStream
);
//cudaCheck(cudaMalloc((void **)&digisF5HB.ids, sizeof(uint32_t) * config.maxChannelsF5HB));
digisF5HB.npresamples = cms::cuda::make_device_unique<uint8_t[]>(
config.maxChannelsF5HB,
cudaStream
);
//cudaCheck(cudaMalloc((void **)&digisF5HB.npresamples, sizeof(uint8_t) * config.maxChannelsF5HB));

digisF3HB.data = cms::cuda::make_device_unique<uint16_t[]>(
config.maxChannelsF3HB * compute_stride<Flavor3>(config.nsamplesF3HB),
cudaStream
);
//cudaCheck(cudaMalloc((void **)&digisF3HB.data,
// config.maxChannelsF3HB * sizeof(uint16_t) * compute_stride<Flavor3>(config.nsamplesF3HB)));
digisF3HB.ids = cms::cuda::make_device_unique<uint32_t[]>(
config.maxChannelsF3HB,
cudaStream
);
//cudaCheck(cudaMalloc((void **)&digisF3HB.ids, config.maxChannelsF3HB * sizeof(uint32_t)));
}
};

struct InputDataGPU {
unsigned char *data = nullptr;
uint32_t *offsets = nullptr;
int *feds = nullptr;

void allocate() {
cudaCheck(cudaMalloc((void **)&data, sizeof(unsigned char) * nbytes_per_fed_max * utca_nfeds_max));
cudaCheck(cudaMalloc((void **)&offsets, sizeof(uint32_t) * utca_nfeds_max));
cudaCheck(cudaMalloc((void **)&feds, sizeof(int) * utca_nfeds_max));
}

void deallocate() {
if (data) {
cudaCheck(cudaFree(data));
cudaCheck(cudaFree(offsets));
cudaCheck(cudaFree(feds));
}
}
cms::cuda::device::unique_ptr<unsigned char[]> data;
cms::cuda::device::unique_ptr<uint32_t[]> offsets;
cms::cuda::device::unique_ptr<int[]> feds;
};

struct ConditionsProducts {
Expand Down
36 changes: 18 additions & 18 deletions EventFilter/HcalRawToDigi/plugins/DecodeGPU.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -568,38 +568,38 @@ namespace hcal {
uint32_t const nbytesTotal) {
// transfer
cudaCheck(cudaMemcpyAsync(
inputGPU.data, inputCPU.data.data(), nbytesTotal * sizeof(unsigned char), cudaMemcpyHostToDevice, cudaStream));
cudaCheck(cudaMemcpyAsync(inputGPU.offsets,
inputCPU.offsets.data(),
inputGPU.data.get(), inputCPU.data.get(), nbytesTotal * sizeof(unsigned char), cudaMemcpyHostToDevice, cudaStream));
cudaCheck(cudaMemcpyAsync(inputGPU.offsets.get(),
inputCPU.offsets.get(),
nfedsWithData * sizeof(uint32_t),
cudaMemcpyHostToDevice,
cudaStream));
cudaCheck(cudaMemsetAsync(scratchGPU.pChannelsCounters, 0, sizeof(uint32_t) * numOutputCollections, cudaStream));
cudaCheck(cudaMemsetAsync(scratchGPU.pChannelsCounters.get(), 0, sizeof(uint32_t) * numOutputCollections, cudaStream));
cudaCheck(cudaMemcpyAsync(
inputGPU.feds, inputCPU.feds.data(), nfedsWithData * sizeof(int), cudaMemcpyHostToDevice, cudaStream));
inputGPU.feds.get(), inputCPU.feds.get(), nfedsWithData * sizeof(int), cudaMemcpyHostToDevice, cudaStream));

// 12 is the max number of modules per crate
kernel_rawdecode_test<32><<<nfedsWithData, 12 * 32, 0, cudaStream>>>(inputGPU.data,
inputGPU.offsets,
inputGPU.feds,
kernel_rawdecode_test<32><<<nfedsWithData, 12 * 32, 0, cudaStream>>>(inputGPU.data.get(),
inputGPU.offsets.get(),
inputGPU.feds.get(),
conditions.eMappingProduct.eid2did,
conditions.eMappingProduct.eid2tid,
outputGPU.digisF01HE.data,
outputGPU.digisF01HE.ids,
outputGPU.digisF5HB.data,
outputGPU.digisF5HB.ids,
outputGPU.digisF5HB.npresamples,
outputGPU.digisF3HB.data,
outputGPU.digisF3HB.ids,
scratchGPU.pChannelsCounters,
outputGPU.digisF01HE.data.get(),
outputGPU.digisF01HE.ids.get(),
outputGPU.digisF5HB.data.get(),
outputGPU.digisF5HB.ids.get(),
outputGPU.digisF5HB.npresamples.get(),
outputGPU.digisF3HB.data.get(),
outputGPU.digisF3HB.ids.get(),
scratchGPU.pChannelsCounters.get(),
config.nsamplesF01HE,
config.nsamplesF5HB,
config.nsamplesF3HB,
nbytesTotal);
cudaCheck(cudaGetLastError());

cudaCheck(cudaMemcpyAsync(outputCPU.nchannels.data(),
scratchGPU.pChannelsCounters,
cudaCheck(cudaMemcpyAsync(outputCPU.nchannels.get(),
scratchGPU.pChannelsCounters.get(),
sizeof(uint32_t) * numOutputCollections,
cudaMemcpyDeviceToHost,
cudaStream));
Expand Down
Loading