Skip to content

Commit

Permalink
Made fission sources and scatter sources array class attributes for S…
Browse files Browse the repository at this point in the history 
…olver classes
  • Loading branch information
wbinventor committed Jul 19, 2013
1 parent c1ceb2e commit 3ca79f5
Show file tree
Hide file tree
Showing 6 changed files with 70 additions and 62 deletions.
32 changes: 16 additions & 16 deletions src/CPUSolver.cpp
View file
Expand Up @@ -217,8 +217,11 @@ void CPUSolver::initializeFluxArrays() {
void CPUSolver::initializeSourceArrays() {

/* Delete old sources arrays if they exist */
if (_fission_source != NULL)
delete [] _fission_source;
if (_fission_sources != NULL)
delete [] _fission_sources;

if (_scatter_sources != NULL)
delete [] _scatter_sources;

if (_source != NULL)
delete [] _source;
Expand All @@ -237,7 +240,8 @@ void CPUSolver::initializeSourceArrays() {
/* Allocate memory for all source arrays */
try{
size = _num_FSRs * _num_groups;
_fission_source = new FP_PRECISION[size];
_fission_sources = new FP_PRECISION[size];
_scatter_sources = new FP_PRECISION[size];
_source = new FP_PRECISION[size];
_old_source = new FP_PRECISION[size];
_reduced_source = new FP_PRECISION[size];
Expand Down Expand Up @@ -500,11 +504,11 @@ void CPUSolver::normalizeFluxes() {
volume = _FSR_volumes[r];

for (int e=0; e < _num_groups; e++)
_fission_source(r,e) = nu_sigma_f[e] * _scalar_flux(r,e) * volume;
_fission_sources(r,e) = nu_sigma_f[e] * _scalar_flux(r,e) * volume;
}

/* Compute the total fission source */
tot_fission_source = pairwise_sum<FP_PRECISION>(_fission_source,
tot_fission_source = pairwise_sum<FP_PRECISION>(_fission_sources,
_num_FSRs*_num_groups);

/* Normalize scalar fluxes in each region */
Expand Down Expand Up @@ -550,6 +554,7 @@ void CPUSolver::normalizeFluxes() {
*/
FP_PRECISION CPUSolver::computeFSRSources() {

int tid;
FP_PRECISION scatter_source;
FP_PRECISION fission_source;
double* nu_sigma_f;
Expand All @@ -565,9 +570,7 @@ FP_PRECISION CPUSolver::computeFSRSources() {
sigma_s, sigma_t, fission_source, scatter_source) schedule(guided)
for (int r=0; r < _num_FSRs; r++) {

FP_PRECISION* scatter_sources = new FP_PRECISION[_num_groups];
FP_PRECISION* fission_sources = new FP_PRECISION[_num_groups];

tid = omp_get_thread_num();
material = _FSR_materials[r];
nu_sigma_f = material->getNuSigmaF();
chi = material->getChi();
Expand All @@ -576,20 +579,20 @@ FP_PRECISION CPUSolver::computeFSRSources() {

/* Compute fission source for each group */
for (int e=0; e < _num_groups; e++)
fission_sources[e] = _scalar_flux(r,e) * nu_sigma_f[e];
_fission_sources(r,e) = _scalar_flux(r,e) * nu_sigma_f[e];

fission_source = pairwise_sum<FP_PRECISION>(fission_sources,
fission_source = pairwise_sum<FP_PRECISION>(&_fission_sources(r,0),
_num_groups);

/* Compute total scattering source for group G */
for (int G=0; G < _num_groups; G++) {
scatter_source = 0;

for (int g=0; g < _num_groups; g++)
scatter_sources[g] = sigma_s[G*_num_groups+g]*_scalar_flux(r,g);
_scatter_sources(r,g) = sigma_s[G*_num_groups+g]*_scalar_flux(r,g);

scatter_source = pairwise_sum<FP_PRECISION>(scatter_sources,
_num_groups);
scatter_source = pairwise_sum<FP_PRECISION>(&_scatter_sources(r,0),
_num_groups);

/* Set the total source for region r in group G */
_source(r,G) = ((1.0 / _k_eff) * fission_source *
Expand All @@ -605,9 +608,6 @@ FP_PRECISION CPUSolver::computeFSRSources() {
/* Update the old source */
_old_source(r,G) = _source(r,G);
}

delete [] scatter_sources;
delete [] fission_sources;
}

/* Sum up the residuals from each group and in each region */
Expand Down
10 changes: 7 additions & 3 deletions src/Solver.cpp
View file
Expand Up @@ -33,7 +33,8 @@ Solver::Solver(Geometry* geometry, TrackGenerator* track_generator) {
_boundary_leakage = NULL;

_scalar_flux = NULL;
_fission_source = NULL;
_fission_sources = NULL;
_scatter_sources = NULL;
_source = NULL;
_old_source = NULL;
_reduced_source = NULL;
Expand Down Expand Up @@ -87,8 +88,11 @@ Solver::~Solver() {
if (_scalar_flux != NULL)
delete [] _scalar_flux;

if (_fission_source != NULL)
delete [] _fission_source;
if (_fission_sources != NULL)
delete [] _fission_sources;

if (_scatter_sources != NULL)
delete [] _scatter_sources;

if (_source != NULL)
delete [] _source;
Expand Down
9 changes: 7 additions & 2 deletions src/Solver.h
View file
Expand Up @@ -31,7 +31,9 @@

#define _boundary_leakage(i,pe2) (_boundary_leakage[2*(i)*_polar_times_groups+(pe2)])

#define _fission_source(r,e) (_fission_source[(r)*_num_groups + (e)])
#define _fission_sources(r,e) (_fission_sources[(r)*_num_groups + (e)])

#define _scatter_sources(r,e) (_scatter_sources[(r)*_num_groups + (e)])

#define _source_residuals(r,e) (_source_residuals[(r)*_num_groups + (e)])

Expand Down Expand Up @@ -120,7 +122,10 @@ class Solver {
FP_PRECISION* _scalar_flux;

/** The fission source in each energy group in each flat source region */
FP_PRECISION* _fission_source;
FP_PRECISION* _fission_sources;

/** The in-scatter source in each energy group in each flat source region */
FP_PRECISION* _scatter_sources;

/** The source in each energy group in each flat source region */
FP_PRECISION* _source;
Expand Down
47 changes: 25 additions & 22 deletions src/VectorizedSolver.cpp
View file
Expand Up @@ -52,9 +52,14 @@ VectorizedSolver::~VectorizedSolver() {
_thread_fsr_flux = NULL;
}

if (_fission_source != NULL) {
_mm_free(_fission_source);
_fission_source = NULL;
if (_fission_sources != NULL) {
_mm_free(_fission_sources);
_fission_sources = NULL;
}

if (_scatter_sources != NULL) {
_mm_free(_scatter_sources);
_scatter_sources = NULL;
}

if (_source != NULL) {
Expand Down Expand Up @@ -175,8 +180,11 @@ void VectorizedSolver::initializeFluxArrays() {
void VectorizedSolver::initializeSourceArrays() {

/* Delete old sources arrays if they exist */
if (_fission_source != NULL)
_mm_free(_fission_source);
if (_fission_sources != NULL)
_mm_free(_fission_sources);

if (_scatter_sources != NULL)
_mm_free(_scatter_sources);

if (_source != NULL)
_mm_free(_source);
Expand All @@ -195,7 +203,8 @@ void VectorizedSolver::initializeSourceArrays() {
/* Allocate aligned memory for all source arrays */
try{
size = _num_FSRs * _num_groups * sizeof(FP_PRECISION);
_fission_source = (FP_PRECISION*)_mm_malloc(size, VEC_ALIGNMENT);
_fission_sources = (FP_PRECISION*)_mm_malloc(size, VEC_ALIGNMENT);
_scatter_sources = (FP_PRECISION*)_mm_malloc(size, VEC_ALIGNMENT);
_source = (FP_PRECISION*)_mm_malloc(size, VEC_ALIGNMENT);
_old_source = (FP_PRECISION*)_mm_malloc(size, VEC_ALIGNMENT);
_reduced_source = (FP_PRECISION*)_mm_malloc(size, VEC_ALIGNMENT);
Expand Down Expand Up @@ -234,18 +243,18 @@ void VectorizedSolver::normalizeFluxes() {
/* Loop over each energy group within this vector */
#pragma simd vectorlength(VEC_LENGTH)
for (int e=v*VEC_LENGTH; e < (v+1)*VEC_LENGTH; e++) {
_fission_source(r,e) = nu_sigma_f[e] * _scalar_flux(r,e);
_fission_source(r,e) *= volume;
_fission_sources(r,e) = nu_sigma_f[e] * _scalar_flux(r,e);
_fission_sources(r,e) *= volume;
}
}
}

/* Compute the total fission source */
int size = _num_FSRs * _num_groups;
#ifdef SINGLE
tot_fission_source = cblas_sasum(size, _fission_source, 1);
tot_fission_source = cblas_sasum(size, _fission_sources, 1);
#else
tot_fission_source = cblas_dasum(size, _fission_source, 1);
tot_fission_source = cblas_dasum(size, _fission_sources, 1);
#endif

/* Compute the normalization factor */
Expand Down Expand Up @@ -304,9 +313,6 @@ FP_PRECISION VectorizedSolver::computeFSRSources() {
sigma_s, sigma_t, fission_source, scatter_source) schedule(guided)
for (int r=0; r < _num_FSRs; r++) {

FP_PRECISION* scatter_sources = new FP_PRECISION[_num_groups];
FP_PRECISION* fission_sources = new FP_PRECISION[_num_groups];

material = _FSR_materials[r];
nu_sigma_f = material->getNuSigmaF();
chi = material->getChi();
Expand All @@ -318,13 +324,13 @@ FP_PRECISION VectorizedSolver::computeFSRSources() {
/* Compute fission source for each group */
#pragma simd vectorlength(VEC_LENGTH)
for (int e=v*VEC_LENGTH; e < (v+1)*VEC_LENGTH; e++)
fission_sources[e] = _scalar_flux(r,e) * nu_sigma_f[e];
_fission_sources(r,e) = _scalar_flux(r,e) * nu_sigma_f[e];
}

#ifdef SINGLE
fission_source = cblas_sasum(_num_groups, fission_sources, 1);
fission_source = cblas_sasum(_num_groups, &_fission_sources(r,0), 1);
#else
fission_source = cblas_dasum(_num_groups, fission_sources, 1);
fission_source = cblas_dasum(_num_groups, &_fission_sources(r,0), 1);
#endif

/* Compute total scattering source for group G */
Expand All @@ -335,13 +341,13 @@ FP_PRECISION VectorizedSolver::computeFSRSources() {

#pragma simd vectorlength(VEC_LENGTH)
for (int g=v*VEC_LENGTH; g < (v+1)*VEC_LENGTH; g++)
scatter_sources[g] = sigma_s[G*_num_groups+g]*_scalar_flux(r,g);
_scatter_sources(r,g) = sigma_s[G*_num_groups+g]*_scalar_flux(r,g);
}

#ifdef SINGLE
scatter_source = cblas_sasum(_num_groups, scatter_sources, 1);
scatter_source = cblas_sasum(_num_groups, &_scatter_sources(r,0), 1);
#else
scatter_source = cblas_dasum(_num_groups, scatter_sources, 1);
scatter_source = cblas_dasum(_num_groups, &_scatter_sources(r,0), 1);
#endif

/* Set the total source for region r in group G */
Expand All @@ -358,9 +364,6 @@ FP_PRECISION VectorizedSolver::computeFSRSources() {
/* Update the old source */
_old_source(r,G) = _source(r,G);
}

delete [] scatter_sources;
delete [] fission_sources;
}

/* Sum up the residuals from each group and in each region */
Expand Down
29 changes: 14 additions & 15 deletions src/dev/gpu/GPUSolver.cu
View file
Expand Up @@ -44,13 +44,13 @@ __constant__ FP_PRECISION inverse_prefactor_spacing[1];
* @param FSR_volumes an array of flat source region volumes
* @param FSR_materials an array of flat source region materials
* @param materials an array of materials on the device
* @param fission_source array of fission sources in each flat source region
* @param fission_sources array of fission sources in each flat source region
*/
__global__ void computeFissionSourcesOnDevice(FP_PRECISION* FSR_volumes,
int* FSR_materials,
dev_material* materials,
FP_PRECISION* scalar_flux,
FP_PRECISION* fission_source) {
FP_PRECISION* fission_sources) {

/* Use a shared memory buffer for each thread's fission source */
extern __shared__ FP_PRECISION shared_fission_source[];
Expand Down Expand Up @@ -85,7 +85,7 @@ __global__ void computeFissionSourcesOnDevice(FP_PRECISION* FSR_volumes,

/* Copy this threads fission source to global memory */
tid = threadIdx.x + blockIdx.x * blockDim.x;
fission_source[tid] = shared_fission_source[threadIdx.x];
fission_sources[tid] = shared_fission_source[threadIdx.x];

return;
}
Expand Down Expand Up @@ -603,7 +603,6 @@ GPUSolver::GPUSolver(Geometry* geom, TrackGenerator* track_generator) :
_materials = NULL;
_dev_tracks = NULL;

_fission_source = NULL;
_tot_absorption = NULL;
_tot_fission = NULL;
_leakage = NULL;
Expand Down Expand Up @@ -678,9 +677,9 @@ GPUSolver::~GPUSolver() {
_FSRs_to_pin_powers = NULL;
}

if (_fission_source != NULL) {
_fission_source_vec.clear();
_fission_source = NULL;
if (_fission_sources != NULL) {
_fission_sources_vec.clear();
_fission_sources = NULL;
}

if (_tot_absorption != NULL) {
Expand Down Expand Up @@ -1185,9 +1184,9 @@ void GPUSolver::initializeThrustVectors() {
log_printf(INFO, "Initializing thrust vectors on the GPU...");

/* Delete old vectors if they exist */
if (_fission_source != NULL) {
_fission_source = NULL;
_fission_source_vec.clear();
if (_fission_sources != NULL) {
_fission_sources = NULL;
_fission_sources_vec.clear();
}

if (_tot_absorption != NULL) {
Expand All @@ -1214,8 +1213,8 @@ void GPUSolver::initializeThrustVectors() {
/* Allocate memory for fission, absorption and source vectors on device */
try{
/* Allocate fission source array on device */
_fission_source_vec.resize(_B * _T);
_fission_source = thrust::raw_pointer_cast(&_fission_source_vec[0]);
_fission_sources_vec.resize(_B * _T);
_fission_sources = thrust::raw_pointer_cast(&_fission_sources_vec[0]);

/* Allocate total absorption reaction rate array on device */
_tot_absorption_vec.resize(_B * _T);
Expand Down Expand Up @@ -1379,10 +1378,10 @@ void GPUSolver::normalizeFluxes() {
_FSR_materials,
_materials,
_scalar_flux,
_fission_source);
_fission_sources);

FP_PRECISION norm_factor = 1.0 / thrust::reduce(_fission_source_vec.begin(),
_fission_source_vec.end());
FP_PRECISION norm_factor = 1.0 / thrust::reduce(_fission_sources_vec.begin(),
_fission_sources_vec.end());

normalizeFluxesOnDevice<<<_B, _T>>>(_scalar_flux, _boundary_flux,
norm_factor);
Expand Down
5 changes: 1 addition & 4 deletions src/dev/gpu/GPUSolver.h
View file
Expand Up @@ -83,9 +83,6 @@ class GPUSolver : public Solver {
/** An array of the cumulative number of tracks for each azimuthal angle */
int* _track_index_offsets;

/** A pointer to the Thrust vector of fission sources in each FSR */
FP_PRECISION* _fission_source;

/** A pointer to the Thrust vector of absorption rates in each FSR */
FP_PRECISION* _tot_absorption;

Expand All @@ -96,7 +93,7 @@ class GPUSolver : public Solver {
FP_PRECISION* _leakage;

/** Thrust vector of fission sources in each FSR */
thrust::device_vector<FP_PRECISION> _fission_source_vec;
thrust::device_vector<FP_PRECISION> _fission_sources_vec;

/** Thrust vector of fission rates in each FSR */
thrust::device_vector<FP_PRECISION> _tot_fission_vec;
Expand Down

0 comments on commit 3ca79f5

Please sign in to comment.