Back to home page

EIC code displayed by LXR
 
 

    

Warning, file /include/Geant4/MCGIDI_vector.hpp was not indexed or was modified since last indexation (in which case cross-reference links may be missing, inaccurate or erroneous).

0001 /*
0002 # <<BEGIN-copyright>>
0003 # Copyright 2019, Lawrence Livermore National Security, LLC.
0004 # This file is part of the gidiplus package (https://github.com/LLNL/gidiplus).
0005 # gidiplus is licensed under the MIT license (see https://opensource.org/licenses/MIT).
0006 # SPDX-License-Identifier: MIT
0007 # <<END-copyright>>
0008 */
0009 
0010 #ifndef MCGIDI_VECTOR_HPP
0011 #define MCGIDI_VECTOR_HPP
0012 
0013 #define CPU_MEM false
0014 #define UVM_MEM true
0015 
0016 #ifdef HAVE_OPENMP_TARGET
0017     #ifdef USE_OPENMP_NO_GPU
0018         #define VAR_MEM false
0019     #else
0020         #define VAR_MEM true
0021     #endif
0022 #else
0023     #define VAR_MEM false
0024 #endif
0025 
0026 typedef int MCGIDI_VectorSizeType;
0027 
0028 #define MCGIDI_SWAP(a,b,type) {type ttttttttt=a;a=b;b=ttttttttt;}
0029 
0030 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0031 #include <cuda.h>
0032 #include <cuda_runtime.h>
0033 #include <cuda_runtime_api.h>
0034 #endif
0035 
0036 #if defined(__HIP__)
0037 #include <hip/hip_version.h>
0038 #include <hip/hip_runtime.h>
0039 #include <hip/hip_runtime_api.h>
0040 #include <hip/hip_common.h>
0041 #endif
0042 
0043 #include <string.h>
0044 #include <stdio.h>
0045 #include "cassert"
0046 #include <algorithm>
0047 #include <LUPI_declareMacro.hpp>
0048 #include <vector>
0049 
0050 namespace MCGIDI {
0051 
0052 template <class T>
0053 class Vector 
0054 {
0055  private:
0056    T* _data;
0057    std::size_t _capacity;
0058    std::size_t _size;
0059    bool _mem_type;
0060 
0061  public:
0062    typedef T* iterator;
0063    typedef T* const_iterator;
0064 
0065    LUPI_HOST_DEVICE Vector()        : _data(0), _capacity(0), _size(0), _mem_type(CPU_MEM) {};
0066    LUPI_HOST_DEVICE Vector( std::size_t s, bool mem_flag = CPU_MEM ) : _data(0), _capacity(s), _size(s), _mem_type(mem_flag)
0067    {
0068        
0069       if( s == 0 ){ _data = nullptr; return;}   
0070         switch ((int)_mem_type){
0071             case CPU_MEM:
0072                 _data = new T [_capacity];
0073                 break;
0074             case UVM_MEM:
0075             {
0076                 void *ptr = nullptr;
0077 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0078                 cudaMallocManaged(&ptr, _capacity*sizeof(T), cudaMemAttachGlobal);
0079 #elif defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)
0080                 hipMallocManaged(&ptr, _capacity*sizeof(T), hipMemAttachGlobal);
0081 #endif
0082                 _data = new(ptr) T[_capacity]; 
0083                 break;
0084             }
0085             default:
0086                 _data = new T [_capacity];
0087                 break;
0088         }
0089    }
0090    LUPI_HOST_DEVICE Vector( std::size_t s, const T& d, bool mem_flag = CPU_MEM ) : _data(0), _capacity(s), _size(s), _mem_type(mem_flag)
0091    { 
0092       if( s == 0 ){ _data = nullptr; return;}   
0093         switch ( (int) _mem_type){
0094             case CPU_MEM:
0095                 _data = new T [_capacity];
0096                 break;
0097             case UVM_MEM:
0098             {
0099                 void *ptr = nullptr;
0100 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0101                 cudaMallocManaged(&ptr, _capacity*sizeof(T), cudaMemAttachGlobal);
0102 #elif defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)
0103                 hipMallocManaged(&ptr, _capacity*sizeof(T), hipMemAttachGlobal);
0104 #endif
0105                 _data = new(ptr) T[_capacity];
0106                 break;
0107             }
0108             default:
0109                 _data = new T [_capacity];
0110                 break;
0111         }
0112       for (std::size_t ii = 0; ii < _capacity; ++ii)
0113          _data[ii] = d;
0114    }
0115 
0116    LUPI_HOST_DEVICE Vector(const Vector<T>& aa )
0117         : _data(0), _capacity(aa._capacity), _size(aa._size), _mem_type(aa._mem_type)
0118    {
0119       if( _capacity == 0 ){ _data = nullptr; return; }
0120 
0121         switch ( (int) _mem_type){
0122             case CPU_MEM:
0123                 _data = new T [_capacity];
0124                 break;
0125             case UVM_MEM:
0126             {
0127                 void *ptr = nullptr;
0128 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0129                 cudaMallocManaged(&ptr, _capacity*sizeof(T), cudaMemAttachGlobal);
0130 #elif defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)
0131                 hipMallocManaged(&ptr, _capacity*sizeof(T), hipMemAttachGlobal);
0132 #endif
0133                 _data = new(ptr) T[_capacity];
0134                 break;
0135             }
0136             default:
0137                 _data = new T [_capacity];
0138                 break;
0139         }
0140  
0141       for (std::size_t ii=0; ii<_size; ++ii)
0142          _data[ii] = aa._data[ii];
0143    }
0144 
0145    LUPI_HOST Vector(const std::vector<T>& aa )
0146         : _data(0), _capacity(aa.size()), _size(aa.size()), _mem_type(CPU_MEM)
0147    {
0148       if( _capacity == 0 ){ _data = nullptr; return;}   
0149 
0150         switch ( (int) _mem_type){
0151             case CPU_MEM:
0152                 _data = new T [_capacity];
0153                 break;
0154             case UVM_MEM:
0155             {
0156                 void *ptr = nullptr;
0157 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0158                 cudaMallocManaged(&ptr, _capacity*sizeof(T), cudaMemAttachGlobal);
0159 #elif defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)
0160                 hipMallocManaged(&ptr, _capacity*sizeof(T), hipMemAttachGlobal);
0161 #endif
0162                 _data = new(ptr) T[_capacity];
0163                 break;
0164             }
0165             default:
0166                 _data = new T [_capacity];
0167                 break;
0168         }
0169  
0170       for (std::size_t ii=0; ii<_size; ++ii)
0171          _data[ii] = aa[ii];
0172    }
0173    
0174    LUPI_HOST_DEVICE ~Vector() { 
0175         switch ( (int) _mem_type){
0176             case CPU_MEM:
0177                 delete[] _data; 
0178                 break;
0179             case UVM_MEM:
0180                  for (std::size_t i=0; i < _size; ++i)
0181                    _data[i].~T();
0182 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0183                 cudaFree(_data);
0184 #elif defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)
0185                 hipFree(_data);
0186 #endif
0187                 break;
0188             default:
0189                 delete[] _data; 
0190                 break;
0191         }
0192    }
0193 
0194    LUPI_HOST_DEVICE iterator begin() { return _data; }
0195 
0196    LUPI_HOST_DEVICE const_iterator begin() const { return _data; }
0197 
0198    LUPI_HOST_DEVICE iterator end() { return _data + _size; }
0199 
0200    LUPI_HOST_DEVICE const_iterator end() const { return _data + _size; }
0201 
0202    /// Needed for copy-swap idiom
0203    LUPI_HOST_DEVICE void swap(Vector<T>& other)
0204    {
0205       MCGIDI_SWAP(_data,     other._data,     T*);
0206       MCGIDI_SWAP(_capacity, other._capacity, std::size_t);
0207       MCGIDI_SWAP(_size,     other._size,     std::size_t);
0208       MCGIDI_SWAP(_mem_type, other._mem_type, bool);
0209    }
0210    
0211    /// Implement assignment using copy-swap idiom
0212    LUPI_HOST_DEVICE Vector<T>& operator=(const Vector<T>& aa)
0213    {
0214       if (&aa != this)
0215       {
0216          Vector<T> temp(aa);
0217          this->swap(temp);
0218       }
0219       return *this;
0220    }
0221 
0222    LUPI_HOST Vector<T>& operator=(const std::vector<T>& aa)
0223    {
0224       Vector<T> temp(aa);
0225       this->swap(temp);
0226       return *this;
0227    }
0228    
0229    LUPI_HOST_DEVICE int get_mem_type()
0230    {
0231     return _mem_type;
0232    }
0233 
0234    LUPI_HOST_DEVICE void push_back( const T& dataElem )
0235    {
0236       assert( _size < _capacity );
0237       _data[_size] = dataElem;
0238       _size++;
0239    }
0240 
0241    LUPI_HOST_DEVICE const T& operator[]( std::size_t index ) const
0242    {
0243       // assert( index < _capacity ); 
0244       // assert( index >= 0); comment out pointless assertion size_t type is >= 0 by definition
0245       return _data[index];
0246    }
0247 
0248    LUPI_HOST_DEVICE T& operator[]( std::size_t index )
0249    {
0250       // assert( index < _capacity );
0251       // assert( index >= 0); comment out pointless assertion size_t type is >= 0 by definition
0252       return _data[index];
0253    }
0254    
0255    LUPI_HOST_DEVICE std::size_t capacity() const
0256    {
0257       return _capacity;
0258    }
0259 
0260    LUPI_HOST_DEVICE std::size_t size() const
0261    {
0262       return _size;
0263    }
0264    
0265    LUPI_HOST_DEVICE T& back()
0266    {
0267       return _data[_size-1];
0268    }
0269    
0270    LUPI_HOST_DEVICE T& back() const
0271    {
0272       return _data[_size-1];
0273    }
0274    
0275    LUPI_HOST_DEVICE void reserve( std::size_t s, char ** address = nullptr, bool mem_flag = CPU_MEM )
0276    {
0277       if (s == _capacity) return;
0278       assert( _capacity == 0 );
0279       _capacity = s;
0280       _mem_type = mem_flag;
0281       if( s == 0 ){ _data = nullptr; return;}   
0282         switch ( (int) _mem_type){
0283             case CPU_MEM:
0284                 if (address == nullptr || *address == nullptr) _data = new T [_capacity];
0285                 else {
0286                     _data = new(*address) T [_capacity];
0287                     *address += sizeof(T) * _capacity;
0288                 }
0289                 break;
0290             case UVM_MEM:
0291             {
0292                 void *ptr = nullptr;
0293 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0294                 cudaMallocManaged(&ptr, _capacity*sizeof(T), cudaMemAttachGlobal);
0295 #elif defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)
0296                 hipMallocManaged(&ptr, _capacity*sizeof(T), hipMemAttachGlobal);
0297 #endif
0298                 _data = new(ptr) T[_capacity];
0299                 break;
0300             }
0301             default:
0302                 if (address == nullptr || *address == nullptr) _data = new T [_capacity];
0303                 else {
0304                     _data = new(*address) T [_capacity];
0305                     *address += sizeof(T) * _capacity;
0306                 }
0307                 break;
0308         }
0309    }
0310 
0311    LUPI_HOST_DEVICE void resize( std::size_t s, char ** address = nullptr, bool mem_flag = CPU_MEM )
0312    {
0313       if (_capacity != 0) { 
0314           assert( _capacity >= s);
0315           _size = s;
0316           return;
0317       }
0318       assert( _capacity == 0 );
0319       _capacity = s;
0320       _size = s;
0321       _mem_type = mem_flag;
0322       if( s == 0 ){ _data = nullptr; return;}   
0323         switch ( (int) _mem_type){
0324             case CPU_MEM:
0325                 if (address == nullptr || *address == nullptr) {
0326                     _data = new T [_capacity];
0327                 }
0328                 else {
0329                     _data = new(*address) T [_capacity];
0330                     std::size_t delta = sizeof(T) * _capacity;
0331                     std::size_t sub = delta % 8;
0332                     if (sub != 0) delta += (8-sub);
0333                     *address += delta;
0334                 }
0335                 break;
0336             case UVM_MEM:
0337             {
0338                 void *ptr = nullptr;
0339 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0340                 cudaMallocManaged(&ptr, _capacity*sizeof(T), cudaMemAttachGlobal);
0341 #elif defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)
0342                 hipMallocManaged(&ptr, _capacity*sizeof(T), hipMemAttachGlobal);
0343 #endif
0344                 _data = new(ptr) T[_capacity];
0345                 break;
0346             }
0347             default:
0348                 if (address == nullptr || *address == nullptr) _data = new T [_capacity];
0349                 else {
0350                     _data = new(*address) T [_capacity];
0351                     std::size_t delta = sizeof(T) * _capacity;
0352                     std::size_t sub = delta % 8;
0353                     if (sub != 0) delta += (8-sub);
0354                     *address += delta;
0355                 }
0356                 break;
0357         }
0358    }
0359 
0360    LUPI_HOST_DEVICE void resize( std::size_t s, const T& d, char ** address = nullptr, bool mem_flag = CPU_MEM ) 
0361    { 
0362       assert( _capacity == 0 );
0363       _capacity = s;
0364       _size = s;
0365       _mem_type = mem_flag;
0366       if( s == 0 ){ _data = nullptr; return;}   
0367         switch ( (int) _mem_type){
0368             case CPU_MEM:
0369                 if (address == nullptr || *address == nullptr) _data = new T [_capacity];
0370                 else {
0371                     _data = new(*address) T [_capacity];
0372                     std::size_t delta = sizeof(T) * _capacity;
0373                     std::size_t sub = delta % 8;
0374                     if (sub != 0) delta += (8-sub);
0375                     *address += delta;
0376                 }
0377                 break;
0378             case UVM_MEM:
0379             {
0380                 void *ptr = nullptr;
0381 #if defined(__CUDACC__) && !defined(__CUDA_ARCH__)
0382                 cudaMallocManaged(&ptr, _capacity*sizeof(T), cudaMemAttachGlobal);
0383 #elif defined(__HIP__) && !defined(__HIP_DEVICE_COMPILE__)
0384                 hipMallocManaged(&ptr, _capacity*sizeof(T), hipMemAttachGlobal);
0385 #endif
0386                 _data = new(ptr) T[_capacity];
0387                 break;
0388             }
0389             default:
0390                 if (address == nullptr || *address == nullptr) _data = new T [_capacity];
0391                 else {
0392                     _data = new(*address) T [_capacity];
0393                     std::size_t delta = sizeof(T) * _capacity;
0394                     std::size_t sub = delta % 8;
0395                     if (sub != 0) delta += (8-sub);
0396                     *address += delta;
0397                     *address += sizeof(T) * _capacity;
0398                 }
0399                 break;
0400         }
0401       for (std::size_t ii = 0; ii < _capacity; ++ii)
0402          _data[ii] = d;
0403    }
0404 
0405    LUPI_HOST_DEVICE bool empty() const
0406    {
0407        return ( _size == 0 );
0408    }
0409 
0410    LUPI_HOST_DEVICE void eraseEnd( std::size_t NewEnd )
0411    {
0412        assert( NewEnd <= _size );
0413        _size = NewEnd;
0414    }
0415 
0416    LUPI_HOST_DEVICE  void pop_back()
0417    {
0418        assert(_size > 0);
0419        _size--;
0420    }
0421 
0422    LUPI_HOST_DEVICE void clear()
0423    {
0424        _size = 0;
0425    }
0426 
0427    LUPI_HOST_DEVICE void appendList( std::size_t listSize, T* list )
0428    {
0429        assert( _size + listSize < _capacity );
0430 
0431        for( std::size_t i = _size; i < _size + listSize; i++ )
0432        {
0433            _data[i] = list[ i-_size ];
0434        }
0435 
0436    }
0437 
0438    //Atomically retrieve an availible index then increment that index some amount
0439    LUPI_HOST_DEVICE std::size_t atomic_Index_Inc( std::size_t inc )
0440    {
0441        if (_size+inc > _capacity)
0442           {MCGIDI_PRINTF("inc too much (size %d, inc %d cap %d)\n", _size, inc, _capacity); abort(); }
0443        assert(_size+inc <= _capacity);
0444        std::size_t pos;
0445 
0446 //       #include "mc_omp_atomic_capture.hh"
0447        {pos = _size; _size = _size + inc;}
0448 
0449        return pos;
0450    }
0451 
0452    // This will not work for a vector of base classes.
0453    LUPI_HOST_DEVICE std::size_t internalSize() const {
0454        std::size_t delta = sizeof(T) * _size;
0455        std::size_t sub = delta % 8;
0456        if (sub != 0) delta += (8-sub);
0457        return delta;
0458    }
0459 
0460    LUPI_HOST_DEVICE void forceCreate(std::size_t a_size, T* a_data) {
0461        _capacity = a_size;
0462        _size = a_size;
0463        _data = a_data;
0464    }
0465 };
0466 
0467 }
0468 #endif
[ Source navigation ] [ Diff markup ] [ Identifier search ] [ general search ]

This page was automatically generated by the 2.3.7 LXR engine.

The LXR team
Valid CSS 2.1! Valid HTML 4.01!