eic-opticks/sysrap/SU.cu

0001 #include "SU.hh"
0002
0003 #include "scuda.h"
0004 #include "squad.h"
0005
0006 #include "sphoton.h"
0007 #include "sphotonlite.h"
0008
0009
0010 #include <thrust/device_ptr.h>
0011 #include <thrust/copy.h>
0012 #include <thrust/count.h>
0013
0014
0015 template<typename T>
0016 T* SU::upload(const T* h, size_t num_items )
0017 {
0018     T* d ;
0019     cudaMalloc(&d, num_items*sizeof(T));
0020     cudaMemcpy(d, h, num_items*sizeof(T), cudaMemcpyHostToDevice);
0021     return d ;
0022 }
0023
0024 template SYSRAP_API quad4*       SU::upload(const quad4* , size_t );
0025 template SYSRAP_API sphoton*     SU::upload(const sphoton* , size_t );
0026 template SYSRAP_API sphotonlite* SU::upload(const sphotonlite* , size_t );
0027
0028
0029 /**
0030 SU::deprecated_select_copy_device_to_host
0031 --------------------------------------------
0032
0033 1. apply thrust::count_if to *d* with *selector* functor yielding *num_select*
0034 2. allocate *d_select* with num_select*sizeof(T) bytes
0035 3. thrust::copy_if from *d* to *d_select* using the *selector* functor
0036 4. host new T[num_select] allocation
0037 5. copies from *d_select* to the *num_select* host array *h* using the selector
0038
0039 This API is deprecated because its awkward as the number selected is not known when making the call.
0040 For example it would be difficult to populate an NP array using this without
0041 making copies.
0042
0043 **/
0044
0045 template<typename T>
0046 void SU::deprecated_select_copy_device_to_host( T** h, unsigned& num_select,  T* d, unsigned num_d, const qselector<T>& selector  )
0047 {
0048     thrust::device_ptr<T> td(d);
0049     num_select = thrust::count_if(td, td+num_d , selector );
0050     std::cout << " num_select " << num_select << std::endl ;
0051
0052     T* d_select ;
0053     cudaMalloc(&d_select,     num_select*sizeof(T));
0054     thrust::device_ptr<T> td_select(d_select);
0055
0056     thrust::copy_if(td, td+num_d , td_select, selector );
0057
0058     *h = new T[num_select] ;
0059     cudaMemcpy(*h, d_select, num_select*sizeof(T), cudaMemcpyDeviceToHost);
0060 }
0061 template SYSRAP_API void     SU::deprecated_select_copy_device_to_host( quad4** h, unsigned& ,  quad4* , unsigned , const qselector<quad4>&  );
0062
0063
0064 /**
0065 SU::count_if
0066 ------------------
0067
0068 1. apply thrust::count_if to *d* with *selector* functor yielding *num_select*
0069
0070 **/
0071
0072 template<typename T>
0073 unsigned SU::count_if( const T* d, unsigned num_d,  const qselector<T>& selector )
0074 {
0075     thrust::device_ptr<const T> td(d);
0076     return thrust::count_if(td, td+num_d , selector );
0077 }
0078
0079 template SYSRAP_API unsigned SU::count_if( const quad4* , unsigned, const qselector<quad4>& );
0080
0081 /**
0082 SU::count_if_sphoton
0083 ----------------------
0084
0085 NB d is device side pointer
0086
0087 **/
0088
0089 size_t SU::count_if_sphoton( const sphoton* d, size_t num_d,  const sphoton_selector& photon_selector )
0090 {
0091     thrust::device_ptr<const sphoton> td(d);
0092     return thrust::count_if(td, td+num_d , photon_selector );
0093 }
0094
0095 size_t SU::count_if_sphotonlite( const sphotonlite* d, size_t num_d,  const sphotonlite_selector& photonlite_selector )
0096 {
0097     thrust::device_ptr<const sphotonlite> td(d);
0098     return thrust::count_if(td, td+num_d , photonlite_selector );
0099 }
0100
0101
0102 /**
0103 SU::device_alloc
0104 -------------------
0105
0106 1. allocates *d* with num*sizeof(T) bytes
0107
0108 **/
0109
0110 template<typename T>
0111 T* SU::device_alloc( unsigned num  )
0112 {
0113     T* d ;
0114     cudaMalloc(&d,  num*sizeof(T));
0115     return d ;
0116 }
0117 template SYSRAP_API char*  SU::device_alloc( unsigned );
0118 template SYSRAP_API float* SU::device_alloc( unsigned );
0119 template SYSRAP_API quad4* SU::device_alloc( unsigned );
0120
0121
0122
0123
0124
0125 template<typename T>
0126 void SU::device_zero( T* d, unsigned num )
0127 {
0128     cudaMemset(d, 0,  num*sizeof(T));
0129 }
0130 template SYSRAP_API void SU::device_zero( quad4*, unsigned );
0131
0132
0133 /**
0134 SU::copy_if_device_to_device_presized
0135 -----------------------------------------
0136
0137 The d_select array must be presized to fit the selection, determine the size using *select_count* with the same selector.
0138
0139 2. thrust::copy_if from *d* to *d_select* using the selector functor
0140 3. copies from *d_select* to the *num_select* presized host array *h* using the selector
0141
0142 **/
0143
0144 template<typename T>
0145 void SU::copy_if_device_to_device_presized( T* d_select, const T* d, unsigned num_d, const qselector<T>& selector )
0146 {
0147     thrust::device_ptr<const T> td(d);
0148     thrust::device_ptr<T> td_select(d_select);
0149     thrust::copy_if(td, td+num_d , td_select, selector );
0150 }
0151
0152 template SYSRAP_API void SU::copy_if_device_to_device_presized( quad4*, const quad4*, unsigned, const qselector<quad4>& );
0153
0154
0155 void SU::copy_if_device_to_device_presized_sphoton( sphoton* d_select, const sphoton* d, size_t num_d, const sphoton_selector& photon_selector )
0156 {
0157     thrust::device_ptr<const sphoton> td(d);
0158     thrust::device_ptr<sphoton> td_select(d_select);
0159     thrust::copy_if(td, td+num_d , td_select, photon_selector );
0160 }
0161
0162 void SU::copy_if_device_to_device_presized_sphotonlite( sphotonlite* d_select, const sphotonlite* d, size_t num_d, const sphotonlite_selector& photonlite_selector )
0163 {
0164     thrust::device_ptr<const sphotonlite> td(d);
0165     thrust::device_ptr<sphotonlite> td_select(d_select);
0166     thrust::copy_if(td, td+num_d , td_select, photonlite_selector );
0167 }
0168
0169
0170
0171 template<typename T>
0172 void SU::copy_device_to_host_presized( T* h, const T* d, unsigned num  )
0173 {
0174     cudaMemcpy(h, d, num*sizeof(T), cudaMemcpyDeviceToHost);
0175 }
0176 template SYSRAP_API void SU::copy_device_to_host_presized( quad4*,   const quad4*,   unsigned );
0177 template SYSRAP_API void SU::copy_device_to_host_presized( sphoton*, const sphoton*, unsigned );
0178 template SYSRAP_API void SU::copy_device_to_host_presized( sphotonlite*, const sphotonlite*, unsigned );
0179
0180
0181
0182
0183
0184 // TODO: error check all these cuda operations following QU
0185
0186
0187 char* SU::device_alloc_sizeof( unsigned num, unsigned sizeof_item )
0188 {
0189     char* d ;
0190     cudaMalloc(&d,  num*sizeof_item );
0191     return d ;
0192 }
0193
0194 void SU::copy_host_to_device_sizeof( char* d, const char* h, unsigned num, unsigned sizeof_item )
0195 {
0196     cudaMemcpy(d, h, num*sizeof_item, cudaMemcpyHostToDevice);
0197 }
0198
0199 void SU::copy_device_to_host_sizeof( char* h, const char* d, unsigned num, unsigned sizeof_item )
0200 {
0201     cudaMemcpy(h, d, num*sizeof_item, cudaMemcpyDeviceToHost);
0202 }
0203
0204 char* SU::upload_array_sizeof(const char* h, unsigned num_items, unsigned sizeof_item ) // static
0205 {
0206     char* d = nullptr ;
0207     cudaMalloc(reinterpret_cast<void**>( &d ), num_items*sizeof_item );
0208     cudaMemcpy(reinterpret_cast<void*>( d ), h, sizeof_item*num_items, cudaMemcpyHostToDevice );
0209     return d ;
0210 }
0211
0212