HcalRecoParamsWithPulseShapesGPU Class Reference

#include <HcalRecoParamsWithPulseShapesGPU.h>

Classes
struct	Product

Public Member Functions
Product const &	getProduct (cudaStream_t) const
	HcalRecoParamsWithPulseShapesGPU (HcalRecoParams const &)
	~HcalRecoParamsWithPulseShapesGPU ()=default

Private Attributes
std::vector< float, cms::cuda::HostAllocator < float > >	acc25nsVec_
std::vector< float, cms::cuda::HostAllocator < float > >	accVarLenIdxMinusOneVec_
std::vector< float, cms::cuda::HostAllocator < float > >	accVarLenIdxZEROVec_
std::vector< float, cms::cuda::HostAllocator < float > >	diff25nsItvlVec_
std::vector< float, cms::cuda::HostAllocator < float > >	diffVarItvlIdxMinusOneVec_
std::vector< float, cms::cuda::HostAllocator < float > >	diffVarItvlIdxZEROVec_
std::vector< uint32_t, cms::cuda::HostAllocator < uint32_t > >	ids_
std::vector< uint32_t, cms::cuda::HostAllocator < uint32_t > >	param1_
std::vector< uint32_t, cms::cuda::HostAllocator < uint32_t > >	param2_
cms::cuda::ESProduct< Product >	product_
uint64_t	totalChannels_

Detailed Description

Definition at line 15 of file HcalRecoParamsWithPulseShapesGPU.h.

Constructor & Destructor Documentation

HcalRecoParamsWithPulseShapesGPU::HcalRecoParamsWithPulseShapesGPU

(

HcalRecoParams const &

recoParams

)

Definition at line 13 of file HcalRecoParamsWithPulseShapesGPU.cc.

References HcalCondObjectContainer< Item >::getAllContainers().

    : totalChannels_{recoParams.getAllContainers()[0].second.size() + recoParams.getAllContainers()[1].second.size()},
      param1_(totalChannels_),
      param2_(totalChannels_),
      ids_(totalChannels_) {
#ifdef HCAL_MAHI_CPUDEBUG
  printf("hello from a reco params with pulse shapes\n");
#endif

  auto const containers = recoParams.getAllContainers();

  HcalPulseShapes pulseShapes;
  std::unordered_map<unsigned int, uint32_t> idCache;

  // fill in eb
  auto const& barrelValues = containers[0].second;
  for (uint64_t i = 0; i < barrelValues.size(); ++i) {
    param1_[i] = barrelValues[i].param1();
    param2_[i] = barrelValues[i].param2();

    auto const pulseShapeId = barrelValues[i].pulseShapeID();
    // FIXME: 0 throws upon look up to HcalPulseShapes
    // although comments state that 0 is reserved,
    // HcalPulseShapes::getShape throws on 0!
    if (pulseShapeId == 0) {
      ids_[i] = 0;
      continue;
    }
    if (auto const iter = idCache.find(pulseShapeId); iter == idCache.end()) {
      // new guy
      auto const newId = idCache.size();
      idCache[pulseShapeId] = newId;
      // this will be the id
      ids_[i] = newId;

      // resize value arrays
      acc25nsVec_.resize(acc25nsVec_.size() + hcal::constants::maxPSshapeBin);
      diff25nsItvlVec_.resize(diff25nsItvlVec_.size() + hcal::constants::maxPSshapeBin);
      accVarLenIdxMinusOneVec_.resize(accVarLenIdxMinusOneVec_.size() + hcal::constants::nsPerBX);
      diffVarItvlIdxMinusOneVec_.resize(diffVarItvlIdxMinusOneVec_.size() + hcal::constants::nsPerBX);
      accVarLenIdxZEROVec_.resize(accVarLenIdxZEROVec_.size() + hcal::constants::nsPerBX);
      diffVarItvlIdxZEROVec_.resize(diffVarItvlIdxZEROVec_.size() + hcal::constants::nsPerBX);

      // precompute and get values from the functor
      auto const& pulseShape = pulseShapes.getShape(pulseShapeId);
      FitterFuncs::PulseShapeFunctor functor{pulseShape, false, false, false, 1, 0, 0, hcal::constants::maxSamples};
      auto const offset256 = newId * hcal::constants::maxPSshapeBin;
      auto const offset25 = newId * hcal::constants::nsPerBX;
      auto const numShapes = newId;
      for (int i = 0; i < hcal::constants::maxPSshapeBin; i++) {
        acc25nsVec_[offset256 * numShapes + i] = functor.acc25nsVec()[i];
        diff25nsItvlVec_[offset256 * numShapes + i] = functor.diff25nsItvlVec()[i];
      }

      for (int i = 0; i < hcal::constants::nsPerBX; i++) {
        accVarLenIdxMinusOneVec_[offset25 * numShapes + i] = functor.accVarLenIdxMinusOneVec()[i];
        diffVarItvlIdxMinusOneVec_[offset25 * numShapes + i] = functor.diffVarItvlIdxMinusOneVec()[i];
        accVarLenIdxZEROVec_[offset25 * numShapes + i] = functor.accVarLenIdxZEROVec()[i];
        diffVarItvlIdxZEROVec_[offset25 * numShapes + i] = functor.diffVarItvlIdxZEROVec()[i];
      }
    } else {
      // already recorded this pulse shape, just set id
      ids_[i] = iter->second;
    }
#ifdef HCAL_MAHI_CPUDEBUG
    if (barrelValues[i].rawId() == DETID_TO_DEBUG) {
      printf("recoShapeId = %u myid = %u\n", pulseShapeId, ids_[i]);
    }
#endif
  }

  // fill in ee
  auto const& endcapValues = containers[1].second;
  auto const offset = barrelValues.size();
  for (uint64_t i = 0; i < endcapValues.size(); ++i) {
    param1_[i + offset] = endcapValues[i].param1();
    param2_[i + offset] = endcapValues[i].param2();

    auto const pulseShapeId = endcapValues[i].pulseShapeID();
    // FIXME: 0 throws upon look up to HcalPulseShapes
    // although comments state that 0 is reserved,
    // HcalPulseShapes::getShape throws on 0!
    if (pulseShapeId == 0) {
      ids_[i + offset] = 0;
      continue;
    }
    if (auto const iter = idCache.find(pulseShapeId); iter == idCache.end()) {
      // new guy
      auto const newId = idCache.size();
      idCache[pulseShapeId] = newId;
      // this will be the id
      ids_[i + offset] = newId;

      // resize value arrays
      acc25nsVec_.resize(acc25nsVec_.size() + hcal::constants::maxPSshapeBin);
      diff25nsItvlVec_.resize(diff25nsItvlVec_.size() + hcal::constants::maxPSshapeBin);
      accVarLenIdxMinusOneVec_.resize(accVarLenIdxMinusOneVec_.size() + hcal::constants::nsPerBX);
      diffVarItvlIdxMinusOneVec_.resize(diffVarItvlIdxMinusOneVec_.size() + hcal::constants::nsPerBX);
      accVarLenIdxZEROVec_.resize(accVarLenIdxZEROVec_.size() + hcal::constants::nsPerBX);
      diffVarItvlIdxZEROVec_.resize(diffVarItvlIdxZEROVec_.size() + hcal::constants::nsPerBX);

      // precompute and get values from the functor
      auto const& pulseShape = pulseShapes.getShape(pulseShapeId);
      FitterFuncs::PulseShapeFunctor functor{pulseShape, false, false, false, 1, 0, 0, hcal::constants::maxSamples};
      auto const offset256 = newId * hcal::constants::maxPSshapeBin;
      auto const offset25 = newId * hcal::constants::nsPerBX;
      auto const numShapes = newId;
      for (int i = 0; i < hcal::constants::maxPSshapeBin; i++) {
        acc25nsVec_[offset256 * numShapes + i] = functor.acc25nsVec()[i];
        diff25nsItvlVec_[offset256 * numShapes + i] = functor.diff25nsItvlVec()[i];
      }

      for (int i = 0; i < hcal::constants::nsPerBX; i++) {
        accVarLenIdxMinusOneVec_[offset25 * numShapes + i] = functor.accVarLenIdxMinusOneVec()[i];
        diffVarItvlIdxMinusOneVec_[offset25 * numShapes + i] = functor.diffVarItvlIdxMinusOneVec()[i];
        accVarLenIdxZEROVec_[offset25 * numShapes + i] = functor.accVarLenIdxZEROVec()[i];
        diffVarItvlIdxZEROVec_[offset25 * numShapes + i] = functor.diffVarItvlIdxZEROVec()[i];
      }
    } else {
      // already recorded this pulse shape, just set id
      ids_[i + offset] = iter->second;
    }
  }

#ifdef HCAL_MAHI_CPUDEBUG
  for (auto const& p : idCache)
    printf("recoPulseShapeId = %u id = %u\n", p.first, p.second);
#endif
}

HcalRecoParamsWithPulseShapesGPU::~HcalRecoParamsWithPulseShapesGPU ( )

default

Member Function Documentation

HcalRecoParamsWithPulseShapesGPU::Product const & HcalRecoParamsWithPulseShapesGPU::getProduct

(

cudaStream_t

cudaStream

)

const

Definition at line 156 of file HcalRecoParamsWithPulseShapesGPU.cc.

References HcalRecoParamsWithPulseShapesGPU::Product::acc25nsVec, acc25nsVec_, HcalRecoParamsWithPulseShapesGPU::Product::accVarLenIdxMinusOneVec, accVarLenIdxMinusOneVec_, HcalRecoParamsWithPulseShapesGPU::Product::accVarLenIdxZEROVec, accVarLenIdxZEROVec_, cudaCheck, HcalRecoParamsWithPulseShapesGPU::Product::diff25nsItvlVec, diff25nsItvlVec_, HcalRecoParamsWithPulseShapesGPU::Product::diffVarItvlIdxMinusOneVec, diffVarItvlIdxMinusOneVec_, HcalRecoParamsWithPulseShapesGPU::Product::diffVarItvlIdxZEROVec, diffVarItvlIdxZEROVec_, HcalRecoParamsWithPulseShapesGPU::Product::ids, ids_, HcalRecoParamsWithPulseShapesGPU::Product::param1, param1_, HcalRecoParamsWithPulseShapesGPU::Product::param2, param2_, and product_.

                                   {
  auto const& product = product_.dataForCurrentDeviceAsync(
      cudaStream, [this](HcalRecoParamsWithPulseShapesGPU::Product& product, cudaStream_t cudaStream) {
        // malloc
        cudaCheck(cudaMalloc((void**)&product.param1, this->param1_.size() * sizeof(uint32_t)));
        cudaCheck(cudaMalloc((void**)&product.param2, this->param2_.size() * sizeof(uint32_t)));
        cudaCheck(cudaMalloc((void**)&product.ids, this->ids_.size() * sizeof(uint32_t)));
        cudaCheck(cudaMalloc((void**)&product.acc25nsVec, this->acc25nsVec_.size() * sizeof(float)));
        cudaCheck(cudaMalloc((void**)&product.diff25nsItvlVec, this->diff25nsItvlVec_.size() * sizeof(float)));
        cudaCheck(cudaMalloc((void**)&product.accVarLenIdxMinusOneVec,
                             this->accVarLenIdxMinusOneVec_.size() * sizeof(float)));
        cudaCheck(cudaMalloc((void**)&product.diffVarItvlIdxMinusOneVec,
                             this->diffVarItvlIdxMinusOneVec_.size() * sizeof(float)));
        cudaCheck(cudaMalloc((void**)&product.accVarLenIdxZEROVec, this->accVarLenIdxZEROVec_.size() * sizeof(float)));
        cudaCheck(
            cudaMalloc((void**)&product.diffVarItvlIdxZEROVec, this->diffVarItvlIdxZEROVec_.size() * sizeof(float)));

        // transfer
        cudaCheck(cudaMemcpyAsync(product.param1,
                                  this->param1_.data(),
                                  this->param1_.size() * sizeof(uint32_t),
                                  cudaMemcpyHostToDevice,
                                  cudaStream));
        cudaCheck(cudaMemcpyAsync(product.param2,
                                  this->param2_.data(),
                                  this->param2_.size() * sizeof(uint32_t),
                                  cudaMemcpyHostToDevice,
                                  cudaStream));
        cudaCheck(cudaMemcpyAsync(
            product.ids, this->ids_.data(), this->ids_.size() * sizeof(uint32_t), cudaMemcpyHostToDevice, cudaStream));
        cudaCheck(cudaMemcpyAsync(product.acc25nsVec,
                                  this->acc25nsVec_.data(),
                                  this->acc25nsVec_.size() * sizeof(float),
                                  cudaMemcpyHostToDevice,
                                  cudaStream));
        cudaCheck(cudaMemcpyAsync(product.diff25nsItvlVec,
                                  this->diff25nsItvlVec_.data(),
                                  this->diff25nsItvlVec_.size() * sizeof(float),
                                  cudaMemcpyHostToDevice,
                                  cudaStream));
        cudaCheck(cudaMemcpyAsync(product.accVarLenIdxMinusOneVec,
                                  this->accVarLenIdxMinusOneVec_.data(),
                                  this->accVarLenIdxMinusOneVec_.size() * sizeof(float),
                                  cudaMemcpyHostToDevice,
                                  cudaStream));
        cudaCheck(cudaMemcpyAsync(product.diffVarItvlIdxMinusOneVec,
                                  this->diffVarItvlIdxMinusOneVec_.data(),
                                  this->diffVarItvlIdxMinusOneVec_.size() * sizeof(float),
                                  cudaMemcpyHostToDevice,
                                  cudaStream));
        cudaCheck(cudaMemcpyAsync(product.accVarLenIdxZEROVec,
                                  this->accVarLenIdxZEROVec_.data(),
                                  this->accVarLenIdxZEROVec_.size() * sizeof(float),
                                  cudaMemcpyHostToDevice,
                                  cudaStream));
        cudaCheck(cudaMemcpyAsync(product.diffVarItvlIdxZEROVec,
                                  this->diffVarItvlIdxZEROVec_.data(),
                                  this->diffVarItvlIdxZEROVec_.size() * sizeof(float),
                                  cudaMemcpyHostToDevice,
                                  cudaStream));
      });

  return product;
}

Member Data Documentation

std::vector<float, cms::cuda::HostAllocator<float> > HcalRecoParamsWithPulseShapesGPU::acc25nsVec_

private

Definition at line 43 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

std::vector<float, cms::cuda::HostAllocator<float> > HcalRecoParamsWithPulseShapesGPU::accVarLenIdxMinusOneVec_

private

Definition at line 45 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

std::vector<float, cms::cuda::HostAllocator<float> > HcalRecoParamsWithPulseShapesGPU::accVarLenIdxZEROVec_

private

Definition at line 47 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

std::vector<float, cms::cuda::HostAllocator<float> > HcalRecoParamsWithPulseShapesGPU::diff25nsItvlVec_

private

Definition at line 44 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

std::vector<float, cms::cuda::HostAllocator<float> > HcalRecoParamsWithPulseShapesGPU::diffVarItvlIdxMinusOneVec_

private

Definition at line 46 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

std::vector<float, cms::cuda::HostAllocator<float> > HcalRecoParamsWithPulseShapesGPU::diffVarItvlIdxZEROVec_

private

Definition at line 48 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

std::vector<uint32_t, cms::cuda::HostAllocator<uint32_t> > HcalRecoParamsWithPulseShapesGPU::ids_

private

Definition at line 41 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

std::vector<uint32_t, cms::cuda::HostAllocator<uint32_t> > HcalRecoParamsWithPulseShapesGPU::param1_

private

Definition at line 39 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

std::vector<uint32_t, cms::cuda::HostAllocator<uint32_t> > HcalRecoParamsWithPulseShapesGPU::param2_

private

Definition at line 40 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

cms::cuda::ESProduct<Product> HcalRecoParamsWithPulseShapesGPU::product_

private

Definition at line 50 of file HcalRecoParamsWithPulseShapesGPU.h.

Referenced by getProduct().

uint64_t HcalRecoParamsWithPulseShapesGPU::totalChannels_

private

Definition at line 38 of file HcalRecoParamsWithPulseShapesGPU.h.