nonbonded_workunit.h

// -*-c++-*-
#ifndef STORMM_NONBONDED_WORKUNIT_H
#define STORMM_NONBONDED_WORKUNIT_H
 
#include "copyright.h"
#include "Accelerator/gpu_details.h"
#include "Math/summation.h"
#include "Topology/atomgraph.h"
#include "Potential/static_exclusionmask.h"
#include "static_mask_synthesis.h"
#include "synthesis_enumerators.h"
 
namespace stormm {
namespace synthesis {
 
using card::GpuDetails;
using energy::StaticExclusionMask;
using energy::tile_length;
using energy::tile_lengths_per_supertile;
using stmath::sum;
  
constexpr int small_block_max_atoms = 320;

constexpr int small_block_max_imports = small_block_max_atoms / tile_length;

constexpr int small_block_tile_width = 8;

constexpr int tiny_nbwu_tiles   =     small_block_tile_width;
constexpr int small_nbwu_tiles  = 2 * small_block_tile_width;
constexpr int medium_nbwu_tiles = 4 * small_block_tile_width;
constexpr int large_nbwu_tiles  = 8 * small_block_tile_width;

constexpr int huge_nbwu_tiles   = tile_lengths_per_supertile * tile_lengths_per_supertile;

constexpr int tile_groups_wu_abstract_length = 64;

constexpr int supertile_wu_abstract_length = 8;

class NonbondedWorkUnit {
public:

  NonbondedWorkUnit(const StaticExclusionMask &se, const std::vector<int3> &tile_list);
  
  NonbondedWorkUnit(const StaticExclusionMask &se, int abscissa_start, int ordinate_start);
 
  NonbondedWorkUnit(const StaticExclusionMaskSynthesis &se,
                    const std::vector<int3> &tile_list);
 
  NonbondedWorkUnit(const StaticExclusionMaskSynthesis &se, int abscissa_start, int ordinate_start,
                    int system_index);

  NonbondedWorkUnit(const NonbondedWorkUnit &original) = default;
  NonbondedWorkUnit(NonbondedWorkUnit &&original) = default;
  NonbondedWorkUnit& operator=(const NonbondedWorkUnit &other) = default;
  NonbondedWorkUnit& operator=(NonbondedWorkUnit &&other) = default;
  
  int getTileCount() const;

  int getImportCount() const;

  int getInitializationMask() const;

  int4 getTileLimits(int index) const;

  const std::vector<uint2>& getTileInstructions() const;
  
  std::vector<int> getAbstract(int instruction_start = 0) const;

  void setInitializationMask(int mask_in);

  void setRefreshAtomIndex(int index_in);
  
  void setRefreshWorkCode(int code_in);
  
private:
  int tile_count;                          
  int import_count;                        
  NbwuKind kind;                           
  int score;                               
  int init_accumulation;                   
  int refresh_atom_start;                  
  int refresh_code;                        
  std::vector<int> imports;                
  std::vector<int> import_system_indices;  
  std::vector<int> import_size_keys;       
  std::vector<uint2> tile_instructions;    

  int getImportSystemIndex(const SeMaskSynthesisReader &ser, int atom_index);
};

bool addTileToWorkUnitList(int3* tile_list, int* import_coverage,
                           const std::vector<int> &system_tile_starts, int *import_count,
                           int *current_tile_count, const int ti, const int tj, const int sysid);

void setPropertyInitializationMasks(std::vector<NonbondedWorkUnit> *all_wu, NbwuKind kind);

int nonbondedWorkUnitInitCode(const InitializationTask init_request, const int cache_depth);

void distributeInitializationRanges(std::vector<NonbondedWorkUnit> *result, int natom,
                                    int init_code, const GpuDetails &gpu = null_gpu);
  
std::vector<NonbondedWorkUnit>
buildNonbondedWorkUnits(const StaticExclusionMaskSynthesis &poly_se,
                        InitializationTask init_request = InitializationTask::NONE,
                        int random_cache_depth = 0, const GpuDetails &gpu = null_gpu);
 
std::vector<NonbondedWorkUnit>
buildNonbondedWorkUnits(const StaticExclusionMask &se,
                        InitializationTask init_request = InitializationTask::NONE,
                        int random_cache_depth = 0);
 
} // namespace synthesis
} // namespace stormm
 
#include "nonbonded_workunit.tpp"
 
#endif