Skip to content

Commit

Permalink
completed (translations between cellX/cellY and physical X/Y missing)…
Browse files Browse the repository at this point in the history 
…: framework for obtaining positions in the GPUs from the detids
  • Loading branch information
@bfonta
Bruno Alves authored and bfonta committed Sep 15, 2020
1 parent 6989f8b commit 4773a0a
Show file tree
Hide file tree
Showing 8 changed files with 174 additions and 70 deletions.
70 changes: 41 additions & 29 deletions CUDADataFormats/HGCal/interface/HGCConditions.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -5,38 +5,42 @@ class HeterogeneousHGCSiliconDetId {
public:
constexpr HeterogeneousHGCSiliconDetId(uint32_t id): id_(id) {}
constexpr uint32_t type() { return (id_ >> kHGCalTypeOffset) & kHGCalTypeMask; }
constexpr uint32_t zside() { return (((id_ >> kHGCalZsideOffset) & kHGCalZsideMask) ? -1 : 1); }
constexpr int32_t zside() { return (((id_ >> kHGCalZsideOffset) & kHGCalZsideMask) ? -1 : 1); }
constexpr uint32_t layer() { return (id_ >> kHGCalLayerOffset) & kHGCalLayerMask; }
constexpr uint32_t waferUAbs() { return (id_ >> kHGCalWaferUOffset) & kHGCalWaferUMask; }
constexpr uint32_t waferVAbs() { return (id_ >> kHGCalWaferVOffset) & kHGCalWaferVMask; }
constexpr uint32_t waferU() { return (((id_ >> kHGCalWaferUSignOffset) & kHGCalWaferUSignMask) ? -waferUAbs() : waferUAbs()); }
constexpr uint32_t waferV() { return (((id_ >> kHGCalWaferVSignOffset) & kHGCalWaferVSignMask) ? -waferVAbs() : waferVAbs()); }
constexpr uint32_t cellU() { return (id_ >> kHGCalCellUOffset) & kHGCalCellUMask; }
constexpr uint32_t cellV() { return (id_ >> kHGCalCellVOffset) & kHGCalCellVMask; }
constexpr int32_t waferUAbs() { return (id_ >> kHGCalWaferUOffset) & kHGCalWaferUMask; }
constexpr int32_t waferVAbs() { return (id_ >> kHGCalWaferVOffset) & kHGCalWaferVMask; }
constexpr int32_t waferU() { return (((id_ >> kHGCalWaferUSignOffset) & kHGCalWaferUSignMask) ? -waferUAbs() : waferUAbs()); }
constexpr int32_t waferV() { return (((id_ >> kHGCalWaferVSignOffset) & kHGCalWaferVSignMask) ? -waferVAbs() : waferVAbs()); }
constexpr int32_t waferX() { return (-2 * waferU() + waferV()); }
constexpr int32_t waferY() { return (2 * waferV()); }
constexpr int32_t cellU() { return (id_ >> kHGCalCellUOffset) & kHGCalCellUMask; }
constexpr int32_t cellV() { return (id_ >> kHGCalCellVOffset) & kHGCalCellVMask; }
constexpr int32_t cellX() { int N = (type() == HGCalFine) ? HGCalFineN : HGCalCoarseN; return (3 * (cellV() - N) + 2); }
constexpr int32_t cellY() { int N = (type() == HGCalFine) ? HGCalFineN : HGCalCoarseN; return (2 * cellU() - (N + cellV())); }

private:
uint32_t id_;
enum waferType { HGCalFine = 0, HGCalCoarseThin = 1, HGCalCoarseThick = 2 };
int HGCalFineN = 12;
int HGCalCoarseN = 8;
int kHGCalCellUOffset = 0;
int kHGCalCellUMask = 0x1F;
int kHGCalCellVOffset = 5;
int kHGCalCellVMask = 0x1F;
int kHGCalWaferUOffset = 10;
int kHGCalWaferUMask = 0xF;
int kHGCalWaferUSignOffset = 14;
int kHGCalWaferUSignMask = 0x1;
int kHGCalWaferVOffset = 15;
int kHGCalWaferVMask = 0xF;
int kHGCalWaferVSignOffset = 19;
int kHGCalWaferVSignMask = 0x1;
int kHGCalLayerOffset = 20;
int kHGCalLayerMask = 0x1F;
int kHGCalZsideOffset = 25;
int kHGCalZsideMask = 0x1;
int kHGCalTypeOffset = 26;
int kHGCalTypeMask = 0x3;
static const int32_t HGCalFineN = 12;
static const int32_t HGCalCoarseN = 8;
static const int32_t kHGCalCellUOffset = 0;
static const int32_t kHGCalCellUMask = 0x1F;
static const int32_t kHGCalCellVOffset = 5;
static const int32_t kHGCalCellVMask = 0x1F;
static const int32_t kHGCalWaferUOffset = 10;
static const int32_t kHGCalWaferUMask = 0xF;
static const int32_t kHGCalWaferUSignOffset = 14;
static const int32_t kHGCalWaferUSignMask = 0x1;
static const int32_t kHGCalWaferVOffset = 15;
static const int32_t kHGCalWaferVMask = 0xF;
static const int32_t kHGCalWaferVSignOffset = 19;
static const int32_t kHGCalWaferVSignMask = 0x1;
static const int32_t kHGCalLayerOffset = 20;
static const int32_t kHGCalLayerMask = 0x1F;
static const int32_t kHGCalZsideOffset = 25;
static const int32_t kHGCalZsideMask = 0x1;
static const int32_t kHGCalTypeOffset = 26;
static const int32_t kHGCalTypeMask = 0x3;
};

class HeterogeneousHGCScintillatorDetId {
Expand Down Expand Up @@ -127,9 +131,12 @@ namespace hgcal_conditions {
float* z;
};

enum class HeterogeneousHGCalPositionsType {Int32_t, Uint32_t};
enum class HeterogeneousHGCalPositionsType {Float, Int32_t, Uint32_t};

const std::vector<HeterogeneousHGCalPositionsType> types = { HeterogeneousHGCalPositionsType::Int32_t,
const std::vector<HeterogeneousHGCalPositionsType> types = { HeterogeneousHGCalPositionsType::Float,
HeterogeneousHGCalPositionsType::Float,
HeterogeneousHGCalPositionsType::Float,
HeterogeneousHGCalPositionsType::Int32_t,
HeterogeneousHGCalPositionsType::Uint32_t };

struct HGCalPositionsMapping {
Expand All @@ -143,6 +150,10 @@ namespace hgcal_conditions {
};

struct HeterogeneousHGCalPositionsMapping {
//the x, y and z positions will not be filled in the CPU
float* x;
float* y;
float* z;
int32_t *numberCellsHexagon;
uint32_t *detid;
//variables required for the mapping of detid -> cell in the geometry
Expand All @@ -160,6 +171,7 @@ namespace hgcal_conditions {
struct HeterogeneousHEFConditionsESProduct {
parameters::HeterogeneousHGCalHEFParameters params;
positions::HeterogeneousHGCalPositionsMapping posmap;
size_t nelems_posmap;
};
struct HeterogeneousHEBConditionsESProduct {
parameters::HeterogeneousHGCalHEBParameters params;
Expand Down
36 changes: 34 additions & 2 deletions RecoLocalCalo/HGCalRecProducers/plugins/HGCalRecHitKernelImpl.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -107,11 +107,12 @@ __global__
void hef_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, const HGChefUncalibratedRecHitConstantData cdata, const hgcal_conditions::HeterogeneousHEFConditionsESProduct* conds, int length)
{
unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;
HeterogeneousHGCSiliconDetId detid(src_soa.id_[tid]);
printf("cellCoarseY: %lf - cellLayer: %d - numberCellsHexagon: %d - DetId: %d - Var: %d\n", conds->params.cellCoarseY_[12], detid.layer(), conds->posmap.numberCellsHexagon[0], conds->posmap.detid[9], conds->posmap.waferMax);

for (unsigned int i = tid; i < length; i += blockDim.x * gridDim.x)
{
HeterogeneousHGCSiliconDetId detid(src_soa.id_[tid]);
printf("cellCoarseY: %lf - cellX: %d - numberCellsHexagon: %d - DetId: %d - Var: %d\n", conds->params.cellCoarseY_[12], detid.cellX(), conds->posmap.numberCellsHexagon[0], conds->posmap.detid[9], conds->posmap.waferMax);

double weight = get_weight_from_layer(detid.layer(), cdata.weights_);
double rcorr = get_thickness_correction(detid.type(), cdata.rcorr_);
double noise = get_noise(detid.type(), cdata.noise_fC_);
Expand All @@ -138,3 +139,34 @@ void heb_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, const
0, 0, 0, 0);
}
}

__global__
void fill_positions_from_detids(const hgcal_conditions::HeterogeneousHEFConditionsESProduct* conds)
{
unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;

for (unsigned int i = tid; i < conds->nelems_posmap; i += blockDim.x * gridDim.x)
{
HeterogeneousHGCSiliconDetId did(conds->posmap.detid[tid]);
uint32_t cU = did.cellU();
uint32_t cV = did.cellV();
uint32_t wU = did.waferU();
uint32_t wV = did.waferV();
conds->posmap.x[tid] = 1.1;
conds->posmap.y[tid] = 1.2;
conds->posmap.z[tid] = 1.3;
}

}

__global__
void print_positions_from_detids(const hgcal_conditions::HeterogeneousHEFConditionsESProduct* conds)
{
unsigned int tid = blockDim.x * blockIdx.x + threadIdx.x;

for (unsigned int i = tid; i < conds->nelems_posmap; i += blockDim.x * gridDim.x)
{
printf("PosX: %lf - PosY: %lf - Posz: %lf\n", conds->posmap.x[tid], conds->posmap.y[tid], conds->posmap.z[tid]);
}

}
8 changes: 8 additions & 0 deletions RecoLocalCalo/HGCalRecProducers/plugins/HGCalRecHitKernelImpl.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,8 @@
#include "CUDADataFormats/HGCal/interface/HGCUncalibratedRecHitsToRecHitsConstants.h"
#include "CUDADataFormats/HGCal/interface/HGCConditions.h"

#include "RecoLocalCalo/HGCalRecProducers/plugins/KernelManagerHGCalRecHit.h"

__global__
void ee_step1(HGCUncalibratedRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGCeeUncalibratedRecHitConstantData cdata, int length);

Expand All @@ -27,4 +29,10 @@ void hef_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGChe
__global__
void heb_to_rechit(HGCRecHitSoA dst_soa, HGCUncalibratedRecHitSoA src_soa, HGChebUncalibratedRecHitConstantData cdata, int length);

__global__
void fill_positions_from_detids(const hgcal_conditions::HeterogeneousHEFConditionsESProduct* conds);

__global__
void print_positions_from_detids(const hgcal_conditions::HeterogeneousHEFConditionsESProduct* conds);

#endif //RecoLocalCalo_HGCalRecProducers_HGCalRecHitKernelImpl_cuh
109 changes: 76 additions & 33 deletions RecoLocalCalo/HGCalRecProducers/plugins/HeterogeneousHGCalHEFConditions.cc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,7 @@ size_t HeterogeneousHGCalHEFConditionsWrapper::allocate_memory_params_(const std
size_t HeterogeneousHGCalHEFConditionsWrapper::allocate_memory_pos_(const std::vector<size_t>& sz)
{
size_t chunk_ = std::accumulate(sz.begin(), sz.end(), 0); //total memory required in bytes
gpuErrchk(cudaMallocHost(&this->posmap_.numberCellsHexagon, chunk_));
gpuErrchk(cudaMallocHost(&this->posmap_.x, chunk_));
return chunk_;
}

Expand Down Expand Up @@ -86,7 +86,9 @@ void HeterogeneousHGCalHEFConditionsWrapper::transfer_data_to_heterogeneous_poin
for(unsigned int i=1; i<cumsum_sizes.size(); ++i) //start at second element (the first is zero)
{
size_t types_size = 0;
if( cpos::types[i-1] == cpos::HeterogeneousHGCalPositionsType::Int32_t )
if( cpos::types[i-1] == cpos::HeterogeneousHGCalPositionsType::Float )
types_size = sizeof(float);
else if( cpos::types[i-1] == cpos::HeterogeneousHGCalPositionsType::Int32_t )
types_size = sizeof(int32_t);
else if( cpos::types[i-1] == cpos::HeterogeneousHGCalPositionsType::Uint32_t )
types_size = sizeof(uint32_t);
Expand All @@ -102,24 +104,35 @@ void HeterogeneousHGCalHEFConditionsWrapper::transfer_data_to_heterogeneous_poin
{
const unsigned int jm1 = j-1;
const size_t shift = cumsum_sizes[j] - cumsum_sizes[jm1];
if( cpos::types[jm1] == cpos::HeterogeneousHGCalPositionsType::Int32_t and
if( cpos::types[jm1] == cpos::HeterogeneousHGCalPositionsType::Float and
cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Float )
select_pointer_f_(&this->posmap_, j) = select_pointer_f_(&this->posmap_, jm1) + shift;
else if( cpos::types[jm1] == cpos::HeterogeneousHGCalPositionsType::Float and
cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Int32_t )
select_pointer_i_(&this->posmap_, j) = reinterpret_cast<int32_t*>( select_pointer_f_(&this->posmap_, jm1) + shift );
else if( cpos::types[jm1] == cpos::HeterogeneousHGCalPositionsType::Int32_t and
cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Uint32_t )
select_pointer_u_(&this->posmap_, j) = reinterpret_cast<uint32_t*>( select_pointer_i_(&this->posmap_, jm1) + shift );
else
edm::LogError("HeterogeneousHGCalHEFConditionsWrapper") << "Wrong HeterogeneousHGCalPositionsMapping type";
}

//copying the pointers' content
for(unsigned int i=cumsum_sizes[j]; i<cumsum_sizes[j+1]; ++i)
if( j>this->number_position_arrays ) //required due to the assymetry between cpos::HeterogeneousHGCalPositionsMapping and cpos::HGCalPositionsMapping
{
unsigned int index = i - cumsum_sizes[j];
if( cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Int32_t ) {
select_pointer_i_(&this->posmap_, j)[index] = select_pointer_i_(cpuPos, j)[index];
}
else if( cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Uint32_t )
for(unsigned int i=cumsum_sizes[j]; i<cumsum_sizes[j+1]; ++i)
{
select_pointer_u_(&this->posmap_, j)[index] = select_pointer_u_(cpuPos, j)[index];
unsigned int index = i - cumsum_sizes[j];
if( cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Int32_t ) {
select_pointer_i_(&this->posmap_, j)[index] = select_pointer_i_(cpuPos, j-this->number_position_arrays)[index];
}
else if( cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Uint32_t )
{
select_pointer_u_(&this->posmap_, j)[index] = select_pointer_u_(cpuPos, j-this->number_position_arrays)[index];
}
else
edm::LogError("HeterogeneousHGCalHEFConditionsWrapper") << "Wrong HeterogeneousHGCalPositions type";
}
else
edm::LogError("HeterogeneousHGCalHEFConditionsWrapper") << "Wrong HeterogeneousHGCalPositions type";
}
}
}
Expand All @@ -129,6 +142,7 @@ void HeterogeneousHGCalHEFConditionsWrapper::transfer_data_to_heterogeneous_vars
this->posmap_.lastLayer = cpuPos->lastLayer;
this->posmap_.waferMin = cpuPos->waferMin;
this->posmap_.waferMax = cpuPos->waferMax;
this->nelems_posmap_ = cpuPos->detid.size();
}

std::vector<size_t> HeterogeneousHGCalHEFConditionsWrapper::calculate_memory_bytes_params_(const HGCalParameters* cpuParams) {
Expand Down Expand Up @@ -162,16 +176,20 @@ std::vector<size_t> HeterogeneousHGCalHEFConditionsWrapper::calculate_memory_byt
std::vector<size_t> sizes(npointers);
for(unsigned int i=0; i<npointers; ++i)
{
if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Int32_t)
sizes[i] = select_pointer_i_(cpuPos, i).size();
if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Float)
sizes[i] = select_pointer_u_(cpuPos, 1).size(); //each position array (x, y and z) will have the same size as the detid array
else if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Int32_t)
sizes[i] = select_pointer_i_(cpuPos, 0).size();
else if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Uint32_t)
sizes[i] = select_pointer_u_(cpuPos, i).size();
sizes[i] = select_pointer_u_(cpuPos, 1).size();
}

std::vector<size_t> sizes_units(npointers);
for(unsigned int i=0; i<npointers; ++i)
{
if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Int32_t)
if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Float)
sizes_units[i] = sizeof(float);
else if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Int32_t)
sizes_units[i] = sizeof(int32_t);
else if(cpos::types[i] == cpos::HeterogeneousHGCalPositionsType::Uint32_t)
sizes_units[i] = sizeof(uint32_t);
Expand Down Expand Up @@ -225,6 +243,22 @@ std::vector<double> HeterogeneousHGCalHEFConditionsWrapper::select_pointer_d_(co
}
}

float*& HeterogeneousHGCalHEFConditionsWrapper::select_pointer_f_(cpos::HeterogeneousHGCalPositionsMapping* cpuObject,
const unsigned int& item) const {
switch(item)
{
case 0:
return cpuObject->x;
case 1:
return cpuObject->y;
case 2:
return cpuObject->z;
default:
edm::LogError("HeterogeneousHGCalHEFConditionsWrapper") << "select_pointer_i(heterogeneous): no item.";
return cpuObject->x;
}
}

int32_t*& HeterogeneousHGCalHEFConditionsWrapper::select_pointer_i_(cpar::HeterogeneousHGCalHEFParameters* cpuObject,
const unsigned int& item) const {
switch(item)
Expand Down Expand Up @@ -253,7 +287,7 @@ int32_t*& HeterogeneousHGCalHEFConditionsWrapper::select_pointer_i_(cpos::Hetero
const unsigned int& item) const {
switch(item)
{
case 0:
case 3:
return cpuObject->numberCellsHexagon;
default:
edm::LogError("HeterogeneousHGCalHEFConditionsWrapper") << "select_pointer_i(heterogeneous): no item.";
Expand All @@ -277,7 +311,7 @@ uint32_t*& HeterogeneousHGCalHEFConditionsWrapper::select_pointer_u_(cpos::Heter
const unsigned int& item) const {
switch(item)
{
case 1:
case 4:
return cpuObject->detid;
default:
edm::LogError("HeterogeneousHGCalHEFConditionsWrapper") << "select_pointer_u(heterogeneous): no item.";
Expand Down Expand Up @@ -311,20 +345,14 @@ hgcal_conditions::HeterogeneousHEFConditionsESProduct const *HeterogeneousHGCalH
gpuErrchk(cudaMallocHost(&data.host, sizeof(hgcal_conditions::HeterogeneousHEFConditionsESProduct)));
// Allocate the payload array(s) on device memory.
gpuErrchk(cudaMalloc(&(data.host->params.cellFineX_), chunk_params_));
gpuErrchk(cudaMalloc(&(data.host->posmap.numberCellsHexagon), chunk_pos_));
gpuErrchk(cudaMalloc(&(data.host->posmap.x), chunk_pos_));
// Complete the host-side information on the payload
data.host->posmap.firstLayer = this->posmap_.firstLayer;
data.host->posmap.lastLayer = this->posmap_.lastLayer;
data.host->posmap.waferMax = this->posmap_.waferMax;
data.host->posmap.waferMin = this->posmap_.waferMin;

// Allocate the payload object on the device memory.
gpuErrchk(cudaMalloc(&data.device, sizeof(hgcal_conditions::HeterogeneousHEFConditionsESProduct)));

// Transfer the payload, first the array(s) ...
gpuErrchk(cudaMemcpyAsync(data.host->params.cellFineX_, this->params_.cellFineX_, chunk_params_, cudaMemcpyHostToDevice, stream));
gpuErrchk(cudaMemcpyAsync(data.host->posmap.numberCellsHexagon, this->posmap_.numberCellsHexagon, chunk_pos_, cudaMemcpyHostToDevice, stream));

data.host->nelems_posmap = this->nelems_posmap_;

//(set the pointers of the parameters)
size_t sdouble = sizeof(double);
for(unsigned int j=0; j<this->sizes_params_.size()-1; ++j)
Expand All @@ -338,19 +366,34 @@ hgcal_conditions::HeterogeneousHEFConditionsESProduct const *HeterogeneousHGCalH
else
edm::LogError("HeterogeneousHGCalHEFConditionsWrapper") << "compare this functions' logic with hgcal_conditions::parameters::typesHEF";
}

//(set the pointers of the positions)

//(set the pointers of the positions' mapping)
size_t sfloat = sizeof(float);
size_t sint32 = sizeof(int32_t);
for(unsigned int j=0; j<this->sizes_pos_.size()-1; ++j)
{
if( cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Int32_t and
cpos::types[j+1] == cpos::HeterogeneousHGCalPositionsType::Uint32_t )
if( cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Float and
cpos::types[j+1] == cpos::HeterogeneousHGCalPositionsType::Float )
select_pointer_f_(&(data.host->posmap), j+1) = select_pointer_f_(&(data.host->posmap), j) + (this->sizes_pos_[j]/sfloat);
else if( cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Float and
cpos::types[j+1] == cpos::HeterogeneousHGCalPositionsType::Int32_t )
select_pointer_i_(&(data.host->posmap), j+1) = reinterpret_cast<int32_t*>( select_pointer_f_(&(data.host->posmap), j) + (this->sizes_pos_[j]/sfloat) );
else if( cpos::types[j] == cpos::HeterogeneousHGCalPositionsType::Int32_t and
cpos::types[j+1] == cpos::HeterogeneousHGCalPositionsType::Uint32_t )
select_pointer_u_(&(data.host->posmap), j+1) = reinterpret_cast<uint32_t*>( select_pointer_i_(&(data.host->posmap), j) + (this->sizes_pos_[j]/sint32) );
}

// Allocate the payload object on the device memory.
gpuErrchk(cudaMalloc(&data.device, sizeof(hgcal_conditions::HeterogeneousHEFConditionsESProduct)));
// Transfer the payload, first the array(s) ...
gpuErrchk(cudaMemcpyAsync(data.host->params.cellFineX_, this->params_.cellFineX_, chunk_params_, cudaMemcpyHostToDevice, stream));

//Important: The transfer does *not* start at posmap.x because the positions are not known in the CPU side!
size_t position_memory_size_to_transfer = chunk_pos_ - this->number_position_arrays*this->nelems_posmap_*sfloat; //size in bytes occupied by the non-position information
gpuErrchk(cudaMemcpyAsync(data.host->posmap.numberCellsHexagon, this->posmap_.numberCellsHexagon, position_memory_size_to_transfer, cudaMemcpyHostToDevice, stream));

// ... and then the payload object
gpuErrchk(cudaMemcpyAsync(data.device, data.host, sizeof(hgcal_conditions::HeterogeneousHEFConditionsESProduct), cudaMemcpyHostToDevice, stream));

}); //gpuData_.dataForCurrentDeviceAsync

// Returns the payload object on the memory of the current device
Expand All @@ -362,7 +405,7 @@ HeterogeneousHGCalHEFConditionsWrapper::GPUData::~GPUData() {
if(host != nullptr)
{
gpuErrchk(cudaFree(host->params.cellFineX_));
gpuErrchk(cudaFree(host->posmap.numberCellsHexagon));
gpuErrchk(cudaFree(host->posmap.x));
gpuErrchk(cudaFreeHost(host));
}
gpuErrchk(cudaFree(device));
Expand Down
Loading

0 comments on commit 4773a0a

Please sign in to comment.