GPUSolver.h

Go to the documentation of this file.

#ifndef GPUSOLVER_H_
#define GPUSOLVER_H_

#ifdef __cplusplus
#ifdef SWIG
#include "Python.h"
#endif
#include "../../Solver.h"
#endif

#define PySys_WriteStdout printf

#include <thrust/copy.h>
#include <iostream>

#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/fill.h>
#include <thrust/reduce.h>
#include <thrust/replace.h>
#include <thrust/functional.h>
#include <thrust/iterator/constant_iterator.h>
#include <thrust/iterator/counting_iterator.h>
#include <thrust/iterator/transform_iterator.h>
#include <thrust/iterator/permutation_iterator.h>
#include <sm_20_atomic_functions.h>
#include "clone.h"
#include "GPUExpEvaluator.h"

#define scalar_flux(tid,e) (scalar_flux[(tid)*(*num_groups) + (e)])

#define old_scalar_flux(tid,e) (old_scalar_flux[(tid)*(*num_groups) + (e)])

#define reduced_sources(tid,e) (reduced_sources[(tid)*(*num_groups) + (e)])

#define fixed_sources(r,e) (fixed_sources[(r)*(*num_groups) + (e)])

#define weights(i,p) (weights[(i)*(*num_polar_2) + (p)])

#define boundary_flux(t,pe2) (boundary_flux[2*(t)*(*polar_times_groups)+(pe2)])

#define start_flux(t,pe2) (start_flux[2*(t)*(*polar_times_groups)+(pe2)])

class GPUSolver : public Solver {

private:

  int _B;

  int _T;

  int* _FSR_materials;

  dev_material* _materials;

  dev_track* _dev_tracks;

  thrust::device_vector<FP_PRECISION> _boundary_flux;

  thrust::device_vector<FP_PRECISION> _start_flux;

  thrust::device_vector<FP_PRECISION> _scalar_flux;

  thrust::device_vector<FP_PRECISION> _old_scalar_flux;

  thrust::device_vector<FP_PRECISION> _fixed_sources;

  thrust::device_vector<FP_PRECISION> _reduced_sources;

  std::map<int, int> _material_IDs_to_indices;

  void copyQuadrature();

public:

  GPUSolver(TrackGenerator* track_generator=NULL);
  virtual ~GPUSolver();

  int getNumThreadBlocks();

  int getNumThreadsPerBlock();
  FP_PRECISION getFSRSource(int fsr_id, int group);
  FP_PRECISION getFlux(int fsr_id, int group);
  void getFluxes(FP_PRECISION* out_fluxes, int num_fluxes);

  void setNumThreadBlocks(int num_blocks);
  void setNumThreadsPerBlock(int num_threads);
  void setGeometry(Geometry* geometry);
  void setTrackGenerator(TrackGenerator* track_generator);
  void setFluxes(FP_PRECISION* in_fluxes, int num_fluxes);

  void initializeExpEvaluator();
  void initializeMaterials(solverMode mode=ADJOINT);
  void initializeFSRs();
  void initializeTracks();
  void initializeFluxArrays();
  void initializeSourceArrays();
  void initializeFixedSources();

  void zeroTrackFluxes();
  void flattenFSRFluxes(FP_PRECISION value);
  void storeFSRFluxes();
  void normalizeFluxes();
  void computeFSRSources();
  void computeFSRFissionSources();
  void computeFSRScatterSources();
  void transportSweep();
  void addSourceToScalarFlux();
  void computeKeff();
  double computeResidual(residualType res_type);

  void computeFSRFissionRates(double* fission_rates, int num_FSRs);
};


#endif /* GPUSOLVER_H_ */