From 228edbe85ad72442eb938f001b683136638e56a8 Mon Sep 17 00:00:00 2001 From: Dan Riley Date: Thu, 10 Jun 2021 16:18:54 -0400 Subject: [PATCH] remove mkfit CUDADataFormat --- .../interface/MkFitSiStripClustersCUDA.h | 144 ------------------ .../src/MkFitSiStripClustersCUDA.cc | 110 ------------- 2 files changed, 254 deletions(-) delete mode 100644 CUDADataFormats/SiStripCluster/interface/MkFitSiStripClustersCUDA.h delete mode 100644 CUDADataFormats/SiStripCluster/src/MkFitSiStripClustersCUDA.cc diff --git a/CUDADataFormats/SiStripCluster/interface/MkFitSiStripClustersCUDA.h b/CUDADataFormats/SiStripCluster/interface/MkFitSiStripClustersCUDA.h deleted file mode 100644 index d9aad6b5ec4fa..0000000000000 --- a/CUDADataFormats/SiStripCluster/interface/MkFitSiStripClustersCUDA.h +++ /dev/null @@ -1,144 +0,0 @@ -#ifndef CUDADataFormats_SiStripCluster_interface_MkFitSiStripClustersCUDA_h -#define CUDADataFormats_SiStripCluster_interface_MkFitSiStripClustersCUDA_h - -#include "CUDADataFormats/SiStripCluster/interface/GPUtypes.h" - -#include "HeterogeneousCore/CUDAUtilities/interface/device_unique_ptr.h" -#include "HeterogeneousCore/CUDAUtilities/interface/host_unique_ptr.h" -#include "HeterogeneousCore/CUDAUtilities/interface/cudaCompat.h" - -#include - -class MkFitSiStripClustersCUDA { -public: - MkFitSiStripClustersCUDA() = default; - explicit MkFitSiStripClustersCUDA(size_t maxClusters, int clustersPerStrip, cudaStream_t stream); - ~MkFitSiStripClustersCUDA() = default; - - MkFitSiStripClustersCUDA(const MkFitSiStripClustersCUDA &) = delete; - MkFitSiStripClustersCUDA &operator=(const MkFitSiStripClustersCUDA &) = delete; - MkFitSiStripClustersCUDA(MkFitSiStripClustersCUDA &&) = default; - MkFitSiStripClustersCUDA &operator=(MkFitSiStripClustersCUDA &&) = default; - - void setNClusters(uint32_t nClusters) { nClusters_h = nClusters; } - - uint32_t nClusters() const { return nClusters_h; } - - class GlobalDeviceView { - public: - // __device__ __forceinline__ float local_xx(int i) const { return __ldg(local_xx_ + i); } - // __device__ __forceinline__ float local_xy(int i) const { return __ldg(local_xy_ + i); } - // __device__ __forceinline__ float local_yy(int i) const { return __ldg(local_yy_ + i); } - // __device__ __forceinline__ float local(int i) const { return __ldg(local_ + i); } - __device__ __forceinline__ float global_x(int i) const { return __ldg(global_x_ + i); } - __device__ __forceinline__ float global_y(int i) const { return __ldg(global_y_ + i); } - __device__ __forceinline__ float global_z(int i) const { return __ldg(global_z_ + i); } - - __device__ __forceinline__ float global_xx(int i) const { return __ldg(global_xx_ + i); } - __device__ __forceinline__ float global_xy(int i) const { return __ldg(global_xy_ + i); } - __device__ __forceinline__ float global_xz(int i) const { return __ldg(global_xz_ + i); } - __device__ __forceinline__ float global_yy(int i) const { return __ldg(global_yy_ + i); } - __device__ __forceinline__ float global_yz(int i) const { return __ldg(global_yz_ + i); } - __device__ __forceinline__ float global_zz(int i) const { return __ldg(global_zz_ + i); } - - __device__ __forceinline__ short layer(int i) const { return __ldg(layer_ + i); } - __device__ __forceinline__ float charge(int i) const { return __ldg(charge_ + i); } - // __device__ __forceinline__ stripgpu::detId_t clusterDetId(int i) const { return __ldg(clusterDetId_ + i); } - // __device__ __forceinline__ uint32_t clusterIndex(int i) const { return __ldg(clusterIndex_ + i); } - // __device__ __forceinline__ uint8_t clusterADCs(int i) const { return __ldg(clusterADCs_ + i); } - __device__ __forceinline__ stripgpu::stripId_t firstStrip(int i) const { return __ldg(firstStrip_ + i); } - __device__ __forceinline__ uint32_t clusterSize(int i) const { return __ldg(clusterSize_ + i); } - - friend MkFitSiStripClustersCUDA; - - // private: - int nClusters_; - - // float *local_xx_; - // float *local_xy_; - // float *local_yy_; - // float *local_; - float *global_x_; - float *global_y_; - float *global_z_; - - float *global_xx_; - float *global_xy_; - float *global_xz_; - float *global_yy_; - float *global_yz_; - float *global_zz_; - - short *layer_; - float *charge_; - stripgpu::detId_t *clusterDetId_; - // uint32_t *clusterIndex_; - //uint8_t *clusterADCs_; - stripgpu::stripId_t *firstStrip_; - uint32_t *clusterSize_; - }; - - GlobalDeviceView *gview() const { return gview_d.get(); } - - class HostView { - public: - explicit HostView(size_t maxClusters, int clustersPerStrip, cudaStream_t stream); - - cms::cuda::host::unique_ptr clusterDetId_h; - // cms::cuda::host::unique_ptr clusterIndex_h; - cms::cuda::host::unique_ptr charge_h; - //cms::cuda::host::unique_ptr clusterADCs_h; - cms::cuda::host::unique_ptr clusterSize_h; - cms::cuda::host::unique_ptr firstStrip_h; - - // cms::cuda::host::unique_ptr local_xx_h; - // cms::cuda::host::unique_ptr local_xy_h; - // cms::cuda::host::unique_ptr local_yy_h; - // cms::cuda::host::unique_ptr local_h; - cms::cuda::host::unique_ptr global_x_h; - cms::cuda::host::unique_ptr global_y_h; - cms::cuda::host::unique_ptr global_z_h; - cms::cuda::host::unique_ptr global_xx_h; - cms::cuda::host::unique_ptr global_xy_h; - cms::cuda::host::unique_ptr global_xz_h; - cms::cuda::host::unique_ptr global_yy_h; - cms::cuda::host::unique_ptr global_yz_h; - cms::cuda::host::unique_ptr global_zz_h; - - cms::cuda::host::unique_ptr layer_h; - int nClusters_h; - }; - - std::unique_ptr hostView(int clustersPerStrip, cudaStream_t stream) const; - -private: - cms::cuda::device::unique_ptr clusterDetId_d; - // cms::cuda::device::unique_ptr clusterIndex_d; - cms::cuda::device::unique_ptr charge_d; - //cms::cuda::device::unique_ptr clusterADCs_d; - cms::cuda::device::unique_ptr clusterSize_d; - cms::cuda::device::unique_ptr firstStrip_d; - - // cms::cuda::device::unique_ptr local_xx_d; - // cms::cuda::device::unique_ptr local_xy_d; - // cms::cuda::device::unique_ptr local_yy_d; - // cms::cuda::device::unique_ptr local_d; - cms::cuda::device::unique_ptr global_x_d; - cms::cuda::device::unique_ptr global_y_d; - cms::cuda::device::unique_ptr global_z_d; - cms::cuda::device::unique_ptr global_xx_d; - cms::cuda::device::unique_ptr global_xy_d; - cms::cuda::device::unique_ptr global_xz_d; - cms::cuda::device::unique_ptr global_yy_d; - cms::cuda::device::unique_ptr global_yz_d; - cms::cuda::device::unique_ptr global_zz_d; - - cms::cuda::device::unique_ptr layer_d; - - cms::cuda::device::unique_ptr gview_d; // "me" pointer - -public: - int nClusters_h; -}; - -#endif diff --git a/CUDADataFormats/SiStripCluster/src/MkFitSiStripClustersCUDA.cc b/CUDADataFormats/SiStripCluster/src/MkFitSiStripClustersCUDA.cc deleted file mode 100644 index 73ebc693a7423..0000000000000 --- a/CUDADataFormats/SiStripCluster/src/MkFitSiStripClustersCUDA.cc +++ /dev/null @@ -1,110 +0,0 @@ -#include "CUDADataFormats/SiStripCluster/interface/MkFitSiStripClustersCUDA.h" -#include "HeterogeneousCore/CUDAUtilities/interface/copyAsync.h" - -MkFitSiStripClustersCUDA::MkFitSiStripClustersCUDA(size_t maxClusters, int clustersPerStrip, cudaStream_t stream) { - clusterDetId_d = cms::cuda::make_device_unique(maxClusters, stream); - // clusterIndex_d = cms::cuda::make_device_unique(maxClusters, stream); - clusterSize_d = cms::cuda::make_device_unique(maxClusters, stream); - charge_d = cms::cuda::make_device_unique(maxClusters, stream); - // clusterADCs_d = cms::cuda::make_device_unique(maxClusters * clustersPerStrip, stream); - firstStrip_d = cms::cuda::make_device_unique(maxClusters, stream); - - // local_xx_d = cms::cuda::make_device_unique(maxClusters, stream); - // local_xy_d = cms::cuda::make_device_unique(maxClusters, stream); - // local_yy_d = cms::cuda::make_device_unique(maxClusters, stream); - // local_d = cms::cuda::make_device_unique(maxClusters, stream); - global_x_d = cms::cuda::make_device_unique(maxClusters, stream); - global_y_d = cms::cuda::make_device_unique(maxClusters, stream); - global_z_d = cms::cuda::make_device_unique(maxClusters, stream); - global_xx_d = cms::cuda::make_device_unique(maxClusters, stream); - global_xy_d = cms::cuda::make_device_unique(maxClusters, stream); - global_xz_d = cms::cuda::make_device_unique(maxClusters, stream); - global_yy_d = cms::cuda::make_device_unique(maxClusters, stream); - global_yz_d = cms::cuda::make_device_unique(maxClusters, stream); - global_zz_d = cms::cuda::make_device_unique(maxClusters, stream); - - layer_d = cms::cuda::make_device_unique(maxClusters, stream); - - auto gview = cms::cuda::make_host_unique(stream); - // gview->local_xx_ = local_xx_d.get(); - // gview->local_xy_ = local_xy_d.get(); - // gview->local_yy_ = local_yy_d.get(); - // gview->local_ = local_d.get(); - gview->global_x_ = global_x_d.get(); - gview->global_y_ = global_y_d.get(); - gview->global_z_ = global_z_d.get(); - gview->global_xx_ = global_xx_d.get(); - gview->global_xy_ = global_xy_d.get(); - gview->global_xz_ = global_xz_d.get(); - gview->global_yy_ = global_yy_d.get(); - gview->global_yz_ = global_yz_d.get(); - gview->global_zz_ = global_zz_d.get(); - gview->charge_ = charge_d.get(); - gview->clusterDetId_ = clusterDetId_d.get(); - // gview->clusterIndex_ = clusterIndex_d.get(); - //gview->clusterADCs_ = clusterADCs_d.get(); - gview->firstStrip_ = firstStrip_d.get(); - gview->clusterSize_ = clusterSize_d.get(); - - gview->layer_ = layer_d.get(); - - gview_d = cms::cuda::make_device_unique(stream); - cms::cuda::copyAsync(gview_d, gview, stream); -} - -MkFitSiStripClustersCUDA::HostView::HostView(size_t maxClusters, int clustersPerStrip, cudaStream_t stream) { - clusterDetId_h = cms::cuda::make_host_unique(maxClusters, stream); - // clusterIndex_h = cms::cuda::make_host_unique(maxClusters, stream); - charge_h = cms::cuda::make_host_unique(maxClusters, stream); - // clusterADCs_h = cms::cuda::make_host_unique(maxClusters * clustersPerStrip, stream); - firstStrip_h = cms::cuda::make_host_unique(maxClusters, stream); - clusterSize_h = cms::cuda::make_host_unique(maxClusters, stream); - - // local_xx_h = cms::cuda::make_host_unique(maxClusters, stream); - // local_xy_h = cms::cuda::make_host_unique(maxClusters, stream); - // local_yy_h = cms::cuda::make_host_unique(maxClusters, stream); - // local_h = cms::cuda::make_host_unique(maxClusters, stream); - global_x_h = cms::cuda::make_host_unique(maxClusters, stream); - global_y_h = cms::cuda::make_host_unique(maxClusters, stream); - global_z_h = cms::cuda::make_host_unique(maxClusters, stream); - global_xx_h = cms::cuda::make_host_unique(maxClusters, stream); - global_xy_h = cms::cuda::make_host_unique(maxClusters, stream); - global_xz_h = cms::cuda::make_host_unique(maxClusters, stream); - global_yy_h = cms::cuda::make_host_unique(maxClusters, stream); - global_yz_h = cms::cuda::make_host_unique(maxClusters, stream); - global_zz_h = cms::cuda::make_host_unique(maxClusters, stream); - - layer_h = cms::cuda::make_host_unique(maxClusters, stream); - - nClusters_h = maxClusters; -} - -std::unique_ptr MkFitSiStripClustersCUDA::hostView(int clustersPerStrip, - cudaStream_t stream) const { - auto view_h = std::make_unique(nClusters_h, clustersPerStrip, stream); - - cms::cuda::copyAsync(view_h->clusterDetId_h, clusterDetId_d, nClusters_h, stream); - // cms::cuda::copyAsync(view_h->clusterIndex_h, clusterIndex_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->charge_h, charge_d, nClusters_h, stream); - // cms::cuda::copyAsync(view_h->clusterADCs_h, clusterADCs_d, nClusters_h * clustersPerStrip, stream); - cms::cuda::copyAsync(view_h->firstStrip_h, firstStrip_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->clusterSize_h, clusterSize_d, nClusters_h, stream); - - // cms::cuda::copyAsync(view_h->local_xx_h, local_xx_d, nClusters_h, stream); - // cms::cuda::copyAsync(view_h->local_xy_h, local_xy_d, nClusters_h, stream); - // cms::cuda::copyAsync(view_h->local_yy_h, local_yy_d, nClusters_h, stream); - // cms::cuda::copyAsync(view_h->local_h, local_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_x_h, global_x_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_y_h, global_y_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_z_h, global_z_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_xx_h, global_xx_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_xy_h, global_xy_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_xz_h, global_xz_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_yy_h, global_yy_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_yz_h, global_yz_d, nClusters_h, stream); - cms::cuda::copyAsync(view_h->global_zz_h, global_zz_d, nClusters_h, stream); - - cms::cuda::copyAsync(view_h->layer_h, layer_d, nClusters_h, stream); - - return view_h; -}