SpatialOps
FieldInfo.h
1 /*
2  * Copyright (c) 2014 The University of Utah
3  *
4  * Permission is hereby granted, free of charge, to any person obtaining a copy
5  * of this software and associated documentation files (the "Software"), to
6  * deal in the Software without restriction, including without limitation the
7  * rights to use, copy, modify, merge, publish, distribute, sublicense, and/or
8  * sell copies of the Software, and to permit persons to whom the Software is
9  * furnished to do so, subject to the following conditions:
10  *
11  * The above copyright notice and this permission notice shall be included in
12  * all copies or substantial portions of the Software.
13  *
14  * THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
15  * IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
16  * FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
17  * AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
18  * LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING
19  * FROM, OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS
20  * IN THE SOFTWARE.
21  * ----------------------------------------------------------------------------
22  * Available debugging flags:
23  *
24  * DEBUG_SF_ALL -- Enable all Spatial Field debugging flags.
25  *
26  */
27 //#define DEBUG_SF_ALL
28 
29 #ifndef SPATIALOPS_FIELDINFO_H
30 #define SPATIALOPS_FIELDINFO_H
31 
32 #include <iostream>
33 #include <stdexcept>
34 #include <sstream>
35 #include <map>
36 
37 #include <spatialops/SpatialOpsConfigure.h>
38 
39 #include <spatialops/structured/ExternalAllocators.h>
40 #include <spatialops/structured/MemoryTypes.h>
45 
46 #ifdef ENABLE_THREADS
47 #include <boost/thread/mutex.hpp>
48 #include <boost/interprocess/sync/scoped_lock.hpp>
49 #endif
50 
51 #ifdef ENABLE_CUDA
52 #include <cuda_runtime.h>
53 #endif
54 
55 namespace SpatialOps{
56 
57  //Forward Declaration
58  template <typename T> class Pool;
59 
66  {
67  InternalStorage,
68  ExternalStorage
69  };
70 
91  template< typename T=double >
92  class FieldInfo {
104  struct DeviceMemory{
105  T* field_;
106  bool isValid_;
107  bool builtField_;
108 
116  DeviceMemory(T* field, bool isValid, bool builtField)
117  : field_(field), isValid_(isValid), builtField_(builtField)
118  {}
119 
127  DeviceMemory()
128  : field_(NULL), isValid_(false), builtField_(false)
129  {}
130  };
131 
132  typedef std::map<short int, DeviceMemory> DeviceMemoryMap;
133  typedef typename DeviceMemoryMap::const_iterator ConstMapIter;
134  typedef typename DeviceMemoryMap::iterator MapIter;
135 
136  DeviceMemoryMap deviceMap_;
137  MemoryWindow wholeWindow_;
138  const BoundaryCellInfo bcInfo_;
139  const GhostData totalGhosts_;
140  GhostData totalValidGhosts_;
141  short int activeDeviceIndex_;
142 
143 # ifdef ENABLE_THREADS
144  int partitionCount_; // Number of partitions Nebo uses in its thread-parallel backend when assigning to this field
145 # endif
146 
147 # ifdef ENABLE_CUDA
148  cudaStream_t cudaStream_;
149  cudaEvent_t cudaEvent_;
150 # endif
151 
152 # ifdef ENABLE_THREADS
153 
157  class ExecMutex {
158  const boost::mutex::scoped_lock lock;
159  inline boost::mutex& get_mutex() const {static boost::mutex m; return m;}
160  public:
161  ExecMutex() : lock( get_mutex() ) {}
162  ~ExecMutex() {}
163  };
164 # endif
165 
166  public:
167  typedef T value_type;
168 
179  FieldInfo( const MemoryWindow& window,
180  const BoundaryCellInfo& bc,
181  const GhostData& ghosts,
182  T* const fieldValues = NULL,
183  const StorageMode mode = InternalStorage,
184  const short int devIdx = CPU_INDEX );
185 
192  ~FieldInfo();
193 
197  const MemoryWindow& window_with_ghost() const { return wholeWindow_; }
198 
204  }
205 
211  get_ghost_data() );
212  }
213 
217  const BoundaryCellInfo& boundary_info() const{ return bcInfo_; }
218 
222  const GhostData& get_ghost_data() const{ return totalGhosts_; }
223 
227  const GhostData& get_valid_ghost_data() const{ return totalValidGhosts_; }
228 
234  inline void reset_valid_ghosts( const GhostData& ghosts ){
235  wholeWindow_ = wholeWindow_.reset_ghosts( totalValidGhosts_,
236  ghosts );
237  totalValidGhosts_ = ghosts;
238  }
239 
243  unsigned int allocated_bytes() const { return sizeof(T) * (wholeWindow_.glob_npts()); }
244 
245 # ifdef ENABLE_THREADS
246 
251  void set_partition_count(const int count) { partitionCount_ = count; }
252 
256  int get_partition_count() const { return partitionCount_; }
257 # endif
258 
259 # ifdef ENABLE_CUDA
260 
265  void set_stream( const cudaStream_t& stream ) { cudaStream_ = stream; }
266 
270  cudaStream_t const & get_stream() const { return cudaStream_; }
271 
275  cudaEvent_t const & get_last_event() const { return cudaEvent_; }
276 # endif
277 
281  inline void wait_for_synchronization() {
282 # ifdef ENABLE_CUDA
283  ema::cuda::CUDADeviceInterface& CDI = ema::cuda::CUDADeviceInterface::self();
284  CDI.sync_stream( get_stream() );
285 # endif
286  }
287 
291  short int active_device_index() const { return activeDeviceIndex_; }
292 
314  void add_device(short int deviceIndex);
315 
337  void add_device_async( short int deviceIndex );
338 
350  void validate_device( short int deviceIndex );
351 
363  void validate_device_async( short int deviceIndex );
364 
375  void set_device_as_active( short int deviceIndex );
376 
377 
388  void set_device_as_active_async( short int deviceIndex );
389 
395  bool is_valid( const short int deviceIndex ) const;
396 
402  bool is_available( const short int deviceIndex ) const;
403 
411  inline T* field_values( const short int deviceIndex = CPU_INDEX );
412 
421  inline const T* const_field_values( const short int deviceIndex = CPU_INDEX ) const;
422  }; // FieldInfo
423 
424 //==================================================================
425 //
426 // Implementation
427 //
428 //==================================================================
429 
430  template<typename T>
432  const BoundaryCellInfo& bc,
433  const GhostData& ghost,
434  T* const fieldValues,
435  const StorageMode mode,
436  const short int devIdx)
437  : wholeWindow_( window ),
438  bcInfo_( bc.limit_by_extent(window.extent()) ),
439  totalGhosts_( ghost.limit_by_extent(window.extent()) ),
440  totalValidGhosts_( ghost.limit_by_extent(window.extent()) ),
441  activeDeviceIndex_( devIdx )
442  //Determine raw byte count -- this is sometimes required for external device allocation.
443 # ifdef ENABLE_THREADS
444  , partitionCount_( get_soft_thread_count() )
445 # endif
446 # ifdef ENABLE_CUDA
447  , cudaStream_( 0 )
448 # endif
449  {
450 # ifdef ENABLE_CUDA
451  cudaEventCreate(&cudaEvent_, cudaEventDisableTiming);
452 # endif
453  //InternalStorage => we build a new field
454  //ExternalStorage => we wrap T*
455 
456  //check if active device index is valid:
457  if( IS_CPU_INDEX(activeDeviceIndex_)
458 # ifdef ENABLE_CUDA
459  || IS_GPU_INDEX(activeDeviceIndex_)
460 # endif
461  ){
462  //set up active device in map:
463  switch( mode ){
464  case InternalStorage:
465  deviceMap_[activeDeviceIndex_] = DeviceMemory( Pool<T>::get( activeDeviceIndex_, window.glob_npts() ), true, true );
466  break;
467  case ExternalStorage:
468  if( fieldValues != NULL || window.local_npts() == 0 ){
469  // allow NULL pointers so long as the window is
470  // empty so that we never dereference the pointer.
471  deviceMap_[activeDeviceIndex_] = DeviceMemory( fieldValues, true, false );
472  }
473  else if( window.local_npts() > 0 ){
474  std::ostringstream msg;
475  msg << "Attempting to use externally allocated memory in FieldInfo constructor, given NULL"
476  << " \n" << "\t - " << __FILE__ << " : " << __LINE__ << std::endl;
477  throw(std::runtime_error(msg.str()));
478  }
479  break;
480  } // switch(mode)
481  }
482  else {
483  //not valid device index
484  std::ostringstream msg;
485  msg << "Attempt to create field on unsupported device ( "
486  << DeviceTypeTools::get_memory_type_description(activeDeviceIndex_)
487  << " )\n" << "\t - " << __FILE__ << " : " << __LINE__ << std::endl;
488  throw(std::runtime_error(msg.str()));
489  }
490  }
491 
492 //------------------------------------------------------------------
493 
494  template<typename T>
496  {
497 # ifdef ENABLE_CUDA
498  cudaEventDestroy(cudaEvent_);
499 # endif
500  for( MapIter iter = deviceMap_.begin(); iter != deviceMap_.end(); ++iter )
501  if( iter->second.builtField_ ) Pool<T>::put(iter->first, iter->second.field_);
502  }
503 
504 //------------------------------------------------------------------
505 
506  template<typename T>
507  void FieldInfo<T>::add_device( const short int deviceIndex )
508  {
509 # ifdef DEBUG_SF_ALL
510  std::cout << "Call to SpatialField::add_device() for device : "
511  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
512 # endif
513 
514  add_device_async(deviceIndex);
516  }
517 
518 //------------------------------------------------------------------
519 
520  template<typename T>
521  void FieldInfo<T>::add_device_async( const short int deviceIndex )
522  {
523 # ifdef DEBUG_SF_ALL
524  std::cout << "Call to SpatialField::add_device_async() for device : "
525  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
526 # endif
527 
528 # ifdef ENABLE_THREADS
529  ExecMutex lock;
530 # endif
531 
532  DeviceTypeTools::check_valid_index(deviceIndex, __FILE__, __LINE__);
533 
534  MapIter iter = deviceMap_.find( deviceIndex );
535 
536  if(iter != deviceMap_.end()) {
537  if(!iter->second.isValid_)
538  validate_device_async( deviceIndex );
539  }
540  else {
541  deviceMap_[deviceIndex] = DeviceMemory(Pool<T>::get( deviceIndex, (allocated_bytes()/sizeof(T)) ),
542  false,
543  true);
544  validate_device_async( deviceIndex );
545  }
546  }
547 
548 //------------------------------------------------------------------
549 
550  template<typename T>
551  void FieldInfo<T>::validate_device( const short int deviceIndex )
552  {
553 # ifdef DEBUG_SF_ALL
554  std::cout << "Call to SpatialField::validate_device() for device : "
555  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
556 # endif
557 
558  validate_device_async(deviceIndex);
560  }
561 
562 //------------------------------------------------------------------
563 
564  template<typename T>
565  void FieldInfo<T>::validate_device_async( const short int deviceIndex )
566  {
567 # ifdef DEBUG_SF_ALL
568  std::cout << "Call to SpatialField::validate_device_async() for device : "
569  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
570 # endif
571 
572  DeviceTypeTools::check_valid_index( deviceIndex, __FILE__, __LINE__ );
573 
574  //check if deviceIndex exists for field
575  MapIter iter = deviceMap_.find( deviceIndex );
576  if(iter == deviceMap_.end()) {
577  std::ostringstream msg;
578  msg << "Error : sync_device() did not find a valid field entry in the map. \n"
579  << "Given: " << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl
580  << "\t - " << __FILE__ << " : " << __LINE__ << std::endl;
581  throw(std::runtime_error(msg.str()));
582  }
583 
584  // check that deviceIndex is not already valid for field
585  if( deviceIndex != active_device_index() &&
586  !deviceMap_[deviceIndex].isValid_ ){
587 
588 # ifdef ENABLE_CUDA
589  ema::cuda::CUDADeviceInterface& CDI = ema::cuda::CUDADeviceInterface::self();
590 # endif
591 
592  if( IS_CPU_INDEX(active_device_index()) &&
593  IS_CPU_INDEX(deviceIndex) ) {
594  // CPU->CPU
595  // should not happen
596  std::ostringstream msg;
597  msg << "Error : sync_device() cannot copy from CPU to CPU did not find a valid field entry in the map. \n"
598  << "\t - " << __FILE__ << " : " << __LINE__ << std::endl;
599  throw( std::runtime_error(msg.str()) );
600  }
601 
602 # ifdef ENABLE_CUDA
603 
604  else if( IS_CPU_INDEX(active_device_index()) &&
605  IS_GPU_INDEX(deviceIndex) ){
606  // CPU->GPU
607 # ifdef DEBUG_SF_ALL
608  std::cout << "data transfer from CPU to GPU" << std::endl;
609 # endif
610  CDI.memcpy_to((void*)deviceMap_[deviceIndex].field_,
611  deviceMap_[active_device_index()].field_,
612  allocated_bytes(),
613  deviceIndex,
614  get_stream());
615  deviceMap_[deviceIndex].isValid_ = true;
616  }
617 
618  else if( IS_GPU_INDEX(active_device_index()) &&
619  IS_CPU_INDEX(deviceIndex) ){
620  //GPU->CPU
621 # ifdef DEBUG_SF_ALL
622  std::cout << "data transfer from GPU to CPU" << std::endl;
623 # endif
624  CDI.memcpy_from((void*)deviceMap_[deviceIndex].field_,
625  deviceMap_[active_device_index()].field_,
626  allocated_bytes(),
628  get_stream());
629  deviceMap_[deviceIndex].isValid_ = true;
630  }
631 
632  else if( IS_GPU_INDEX(active_device_index()) &&
633  IS_GPU_INDEX(deviceIndex) ){
634  //GPU->GPU
635 # ifdef DEBUG_SF_ALL
636  std::cout << "data transfer from GPU to GPU" << std::endl;
637 # endif
638  CDI.memcpy_peer((void*)deviceMap_[deviceIndex].field_,
639  deviceIndex,
640  deviceMap_[active_device_index()].field_,
642  allocated_bytes());
643  deviceMap_[deviceIndex].isValid_ = true;
644  }
645 # endif // ENABLE_CUDA
646 
647  else{
648  std::ostringstream msg;
649  msg << "Error : sync_device() called on the field with: "
650  << DeviceTypeTools::get_memory_type_description(deviceIndex)
651  << "\n\t - " << __FILE__ << " : " << __LINE__ << std::endl;
652  throw(std::runtime_error(msg.str()));
653  }
654  }
655  }
656 
657 //------------------------------------------------------------------
658 
659  template<typename T>
660  void FieldInfo<T>::set_device_as_active( const short int deviceIndex )
661  {
662 # ifdef DEBUG_SF_ALL
663  std::cout << "Call to SpatialField::set_device_as_active() for device : "
664  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
665 # endif
666 
667  set_device_as_active_async( deviceIndex );
669  }
670 
671 //------------------------------------------------------------------
672 
673  template<typename T>
674  void FieldInfo<T>::set_device_as_active_async( const short int deviceIndex )
675  {
676 # ifdef DEBUG_SF_ALL
677  std::cout << "Call to SpatialField::set_device_as_active() for device : "
678  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
679 # endif
680 
681  DeviceTypeTools::check_valid_index( deviceIndex, __FILE__, __LINE__ );
682 
683  // make device available and valid, if not already
684  if( deviceMap_.find(deviceIndex) == deviceMap_.end() ||
685  !deviceMap_[deviceIndex].isValid_ ){
686  add_device_async( deviceIndex );
687  }
688  activeDeviceIndex_ = deviceIndex;
689  }
690 
691 //------------------------------------------------------------------
692 
693  template<typename T>
694  bool FieldInfo<T>::is_valid( const short int deviceIndex ) const
695  {
696 # ifdef DEBUG_SF_ALL
697  std::cout << "Call to SpatialField::is_valid() for device : "
698  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
699 # endif
700 
701  DeviceTypeTools::check_valid_index( deviceIndex, __FILE__, __LINE__ );
702 
703  ConstMapIter iter = deviceMap_.find( deviceIndex );
704 # ifndef DEBUG_SF_ALL
705  if( iter == deviceMap_.end() )
706  std::cout << "Field Location " << DeviceTypeTools::get_memory_type_description( deviceIndex )
707  << " is not allocated. " << std::endl;
708  else if( !iter->second.isValid_ )
709  std::cout << "Field Location " << DeviceTypeTools::get_memory_type_description( deviceIndex )
710  << " is not valid. " << std::endl;
711 # endif
712  return ( iter != deviceMap_.end() && iter->second.isValid_ );
713  }
714 
715 //------------------------------------------------------------------
716 
717  template<typename T>
718  bool FieldInfo<T>::is_available( const short int deviceIndex ) const
719  {
720 # ifdef DEBUG_SF_ALL
721  std::cout << "Call to SpatialField::is_available() for device : "
722  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
723 # endif
724 
725  DeviceTypeTools::check_valid_index( deviceIndex, __FILE__, __LINE__ );
726 
727  ConstMapIter iter = deviceMap_.find( deviceIndex );
728 # ifndef DEBUG_SF_ALL
729  if( iter == deviceMap_.end() )
730  std::cout << "Field Location " << DeviceTypeTools::get_memory_type_description( deviceIndex )
731  << " is not allocated. " << std::endl;
732 # endif
733  return ( iter != deviceMap_.end() );
734  }
735 
736 //------------------------------------------------------------------
737 
738  template<typename T>
739  T* FieldInfo<T>::field_values( const short int deviceIndex )
740  {
741 # ifdef DEBUG_SF_ALL
742  std::cout << "Call to non-const SpatialField::field_values() for device : "
743  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
744 # endif
745 
746  DeviceTypeTools::check_valid_index( deviceIndex, __FILE__, __LINE__ );
747 
748  // check active
749  if( active_device_index() == deviceIndex ){
750  MapIter iter = deviceMap_.find( deviceIndex );
751 
752  // mark all the other devices except active device as invalid
753  for( MapIter iter2 = deviceMap_.begin(); iter2 != deviceMap_.end(); ++iter2 ){
754  if( !(iter2->first == activeDeviceIndex_) ) iter2->second.isValid_ = false;
755  }
756  return iter->second.field_;
757  }
758  else{
759  std::ostringstream msg;
760  msg << "Request for nonconst field pointer on a nonactive device: "
761  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl
762  << "Active device is: "
763  << DeviceTypeTools::get_memory_type_description(active_device_index()) << std::endl
764  << "Please check the arguments passed into the function."
765  << "\t - " << __FILE__ << " : " << __LINE__ << std::endl;
766  throw( std::runtime_error(msg.str()) );
767  }
768  }
769 
770 //------------------------------------------------------------------
771 
772  template<typename T>
773  const T* FieldInfo<T>::const_field_values( const short int deviceIndex ) const
774  {
775 # ifdef DEBUG_SF_ALL
776  std::cout << "Call to const SpatialField::field_values() for device : "
777  << DeviceTypeTools::get_memory_type_description(deviceIndex) << std::endl;
778 # endif
779 
780  DeviceTypeTools::check_valid_index( deviceIndex, __FILE__, __LINE__ );
781 
782  ConstMapIter iter = deviceMap_.find( deviceIndex );
783 
784  if( iter == deviceMap_.end() ) {
785  //device is not available
786  std::ostringstream msg;
787  msg << "\nRequest for const field pointer on a device for which it has not been allocated\n"
788  << DeviceTypeTools::get_memory_type_description(deviceIndex)
789  << "\t - " << __FILE__ << " : " << __LINE__ << std::endl;
790 # ifndef __CUDACC__
791  msg << "\nThis may be caused by not compiling your source with NVCC."
792  << "\nYou should check the build system.\n\n";
793 # endif
794  throw(std::runtime_error(msg.str()));
795  }
796 // else {
797 // //device is available but not valid
798 // std::ostringstream msg;
799 // msg << "Requested const field pointer on a device is not valid! \n"
800 // << "\t - " << __FILE__ << " : " << __LINE__ << std::endl;
801 // throw(std::runtime_error(msg.str()));
802 // }
803  return iter->second.field_;
804  }
805 
806 //------------------------------------------------------------------
807 
808 } // namespace SpatialOps
809 
810 #endif // SPATIALOPS_FIELDINFO_H
T * field_values(const short int deviceIndex=CPU_INDEX)
return a non-constant pointer to memory on the given device
Definition: FieldInfo.h:739
MemoryWindow remove_ghosts(GhostData const &ghosts) const
return the current memory window with given ghosts removed
Definition: MemoryWindow.h:214
Provides tools to index into a sub-block of memory.
Definition: MemoryWindow.h:63
void add_device_async(short int deviceIndex)
add device memory to this field for given device and populate it with values from current active devi...
Definition: FieldInfo.h:521
void reset_valid_ghosts(const GhostData &ghosts)
set the number of valid ghost cells to given ghosts
Definition: FieldInfo.h:234
void set_device_as_active_async(short int deviceIndex)
set given device (deviceIndex) as active ASYNCHRONOUS VERSION
Definition: FieldInfo.h:674
const T * const_field_values(const short int deviceIndex=CPU_INDEX) const
return a constant pointer to memory on the given device
Definition: FieldInfo.h:773
const MemoryWindow window_without_ghost() const
return the global memory window with no ghost cells
Definition: FieldInfo.h:202
static void put(const short int deviceLocation, T *t)
Return the requested block of memory to the pool. This should only be done for memory requested from ...
Definition: MemoryPool.cpp:217
const GhostData & get_valid_ghost_data() const
return the number of current valid ghost cells
Definition: FieldInfo.h:227
~FieldInfo()
FieldInfo destructor.
Definition: FieldInfo.h:495
void add_device(short int deviceIndex)
add device memory to this field for given device and populate it with values from current active devi...
Definition: FieldInfo.h:507
const MemoryWindow & window_with_ghost() const
return the global memory window with valid ghost cells
Definition: FieldInfo.h:197
size_t glob_npts() const
obtain the global number of points in the field. Note that this is not necessarily contiguous memory ...
Definition: MemoryWindow.h:155
MemoryWindow reset_ghosts(GhostData const &oldGhosts, GhostData const &newGhosts) const
return the current memory window with oldGhosts removed and newGhosts added
Definition: MemoryWindow.h:198
const GhostData & get_ghost_data() const
return the number of total ghost cells
Definition: FieldInfo.h:222
const BoundaryCellInfo & boundary_info() const
return the boundary cell information
Definition: FieldInfo.h:217
size_t local_npts() const
obtain the local number of points in the field. Note that this is not necessarily contiguous memory ...
Definition: MemoryWindow.h:161
bool is_valid(const short int deviceIndex) const
check if the device (deviceIndex) is available and valid
Definition: FieldInfo.h:694
void validate_device_async(short int deviceIndex)
popluate the memory of the given device (deviceIndex) with values from the active device ASYNCHRONOUS...
Definition: FieldInfo.h:565
StorageMode
Specifies how memory should be treated in a SpatialField.
Definition: FieldInfo.h:65
Holds information about the number of ghost cells on each side of the domain.
Definition: GhostData.h:54
bool is_available(const short int deviceIndex) const
check if the device (deviceIndex) is available
Definition: FieldInfo.h:718
void validate_device(short int deviceIndex)
popluate the memory of the given device (deviceIndex) with values from the active device SYNCHRONOUS ...
Definition: FieldInfo.h:551
short int active_device_index() const
return the index of the current active device
Definition: FieldInfo.h:291
const MemoryWindow window_with_all_ghost() const
return the global memory window with all possible ghost cells
Definition: FieldInfo.h:209
unsigned int allocated_bytes() const
returns the number of allocated bytes (for copying memory)
Definition: FieldInfo.h:243
FieldInfo(const MemoryWindow &window, const BoundaryCellInfo &bc, const GhostData &ghosts, T *const fieldValues=NULL, const StorageMode mode=InternalStorage, const short int devIdx=CPU_INDEX)
FieldInfo constructor.
Definition: FieldInfo.h:431
void set_device_as_active(short int deviceIndex)
set given device (deviceIndex) as active SYNCHRONOUS VERSION
Definition: FieldInfo.h:660
Abstracts the internals of a field.
Definition: FieldInfo.h:92
Provides information about boundary cells for various fields.
void wait_for_synchronization()
wait until the current stream is done with all work
Definition: FieldInfo.h:281