diff --git a/src/CPUSolver.cpp b/src/CPUSolver.cpp index 4bcfee8fc..09e6ec770 100644 --- a/src/CPUSolver.cpp +++ b/src/CPUSolver.cpp @@ -413,7 +413,7 @@ void CPUSolver::initializeFSRs() { _num_fissionable_FSRs++; log_printf(DEBUG, "FSR ID = %d has Material ID = %d " - "and volume = %f", r, _FSR_materials[r]->getUid(), + "and volume = %f", r, _FSR_materials[r]->getId(), _FSR_volumes[r]); } diff --git a/src/Geometry.cpp b/src/Geometry.cpp index 44ae69026..de68b6ecd 100644 --- a/src/Geometry.cpp +++ b/src/Geometry.cpp @@ -739,26 +739,19 @@ void Geometry::subdivideCells() { /** * @brief Compute the number of flat source regions in the Geometry and - * initialize arrays for FSR IDs and maps. + * initialize CMFD. * @details This method is intended to be called by the user before initiating * source iteration. This method first subdivides all Cells by calling - * the Geometry::subdivideCells() method. Then it computes the total - * number of FSRs in the Geometry and initializes integer arrays of - * maps to Cells and Materials UIDs/IDs indexed by FSR IDs. + * the Geometry::subdivideCells() method. Then it initializes the CMFD + * object. */ void Geometry::initializeFlatSourceRegions() { /* Subdivide Cells into sectors and rings */ subdivideCells(); - /* Assign UIDs to materials */ + /* Create map of Material IDs to Material pointers */ _all_materials = getAllMaterials(); - std::map::iterator iter; - int uid = 0; - for (iter = _all_materials.begin(); iter != _all_materials.end(); ++iter){ - iter->second->setUid(uid); - uid++; - } /* Initialize CMFD */ if (_cmfd != NULL) diff --git a/src/Material.cpp b/src/Material.cpp index 2b4f80347..5f092238d 100644 --- a/src/Material.cpp +++ b/src/Material.cpp @@ -42,8 +42,6 @@ Material::Material(int id, const char* name) { else _id = id; - _uid = -1; - _name = NULL; setName(name); @@ -145,23 +143,6 @@ Material::~Material() { } -/** - * @brief Set the Material's unique ID. - * @param uid the Material's unique ID - */ -void Material::setUid(int uid) { - _uid = uid; -} - - -/** - * @brief Return the Material's unique ID. - * @return the Material's unique ID - */ -int Material::getUid() const { - return _uid; -} - /** * @brief Return the Material's user-defined ID * @return the Material's user-defined ID @@ -397,7 +378,7 @@ FP_PRECISION Material::getSigmaSByGroup(int origin, int destination) { if (origin <= 0 || destination <= 0 || origin > _num_groups || destination > _num_groups) log_printf(ERROR, "Unable to get sigma_s for group %d,%d for Material %d " "which contains %d energy groups", - origin, destination, _uid, _num_groups); + origin, destination, _id, _num_groups); return getSigmaSByGroupInline(origin-1,destination-1); diff --git a/src/Material.h b/src/Material.h index a16d9960a..458bb200b 100644 --- a/src/Material.h +++ b/src/Material.h @@ -53,9 +53,6 @@ class Material { private: - /** A monotonically increasing unique ID for each Material created */ - int _uid; - /** A user-defined ID for each Material created */ int _id; @@ -114,8 +111,6 @@ class Material { Material(int id=0, const char* name=""); virtual ~Material(); - void setUid(int uid); - int getUid() const; int getId() const; char* getName() const; int getNumEnergyGroups() const; diff --git a/src/accel/DeviceMaterial.h b/src/accel/DeviceMaterial.h index 86f6d596e..660bac719 100644 --- a/src/accel/DeviceMaterial.h +++ b/src/accel/DeviceMaterial.h @@ -21,9 +21,6 @@ */ struct dev_material { - /** A monotonically increasing unique ID for each Material created */ - int _uid; - /** A user-defined ID for each Material created */ int _id; diff --git a/src/accel/DeviceTrack.h b/src/accel/DeviceTrack.h index 7e9818aab..beec94e44 100644 --- a/src/accel/DeviceTrack.h +++ b/src/accel/DeviceTrack.h @@ -25,8 +25,8 @@ struct dev_segment { /** The length of the segment (cm) */ FP_PRECISION _length; - /** A pointer to the Material in which this segment resides */ - int _material_uid; + /** An index into the _materials array that contains Material pointers */ + int _material_index; /** The ID for flat source region in which this segment resides */ int _region_uid; diff --git a/src/accel/cuda/GPUSolver.cu b/src/accel/cuda/GPUSolver.cu index d5e5dbea0..9921ed81a 100644 --- a/src/accel/cuda/GPUSolver.cu +++ b/src/accel/cuda/GPUSolver.cu @@ -23,6 +23,7 @@ __constant__ FP_PRECISION sinthetas[MAX_POLAR_ANGLES]; /** An array of the weights for the polar angles from the Quadrature set */ __constant__ FP_PRECISION polar_weights[MAX_POLAR_ANGLES*MAX_AZIM_ANGLES]; + /** A pointer to an array with the number of tracks per azimuthal angle */ __constant__ int num_tracks[MAX_AZIM_ANGLES/2]; @@ -69,7 +70,7 @@ __device__ int round_to_int(double x) { /** * @brief Compute the total fission source from all FSRs on the GPU. * @param FSR_volumes an array of FSR volumes - * @param FSR_materials an array of FSR Material UIDs + * @param FSR_materials an array of FSR Material indices * @param materials an array of dev_materials on the device * @param scalar_flux the scalar flux in each FSR and energy group * @param fission_sources array of fission sources in each FSR and energy group @@ -167,7 +168,7 @@ __global__ void normalizeFluxesOnDevice(FP_PRECISION* scalar_flux, * \f$ res = \sqrt{\frac{\displaystyle\sum \displaystyle\sum * \left(\frac{Q^i - Q^{i-1}{Q^i}\right)^2}{# FSRs}}} $\f * - * @param FSR_materials an array of FSR Material UIDs + * @param FSR_materials an array of FSR Material indices * @param materials an array of dev_material pointers * @param scalar_flux an array of FSR scalar fluxes * @param old_fission_sources an array of current FSR sources from previous iteration @@ -255,7 +256,7 @@ __global__ void computeFSRSourcesOnDevice(int* FSR_materials, * @brief Compute the total fission source from all FSRs and energy groups * on the GPU. * @param FSR_volumes an array of the FSR volumes - * @param FSR_materials an array of the FSR Material UIDs + * @param FSR_materials an array of the FSR Material indices * @param materials an array of the dev_material pointers * @param scalar_flux an array of FSR scalar fluxes * @param total array of FSR total reaction rates @@ -423,7 +424,7 @@ __device__ void scalarFluxTally(dev_segment* curr_segment, int fsr_id = curr_segment->_region_uid; FP_PRECISION length = curr_segment->_length; - dev_material* curr_material = &materials[curr_segment->_material_uid]; + dev_material* curr_material = &materials[curr_segment->_material_index]; FP_PRECISION *sigma_t = curr_material->_sigma_t; /* The change in angular flux long this Track segment in this FSR */ @@ -616,7 +617,7 @@ __global__ void transportSweepOnDevice(FP_PRECISION* scalar_flux, * @param scalar_flux an array of FSR scalar fluxes * @param reduced_sources an array of FSR sources / total xs * @param FSR_volumes an array of FSR volumes - * @param FSR_materials an array of FSR material UIDs + * @param FSR_materials an array of FSR material indices * @param materials an array of dev_material pointers */ __global__ void addSourceToScalarFluxOnDevice(FP_PRECISION* scalar_flux, @@ -660,7 +661,7 @@ __global__ void addSourceToScalarFluxOnDevice(FP_PRECISION* scalar_flux, * GPUSolver::computeFSRFissionRates(...) method. * @param fission_rates an array to store the fission rates * @param fission_rates an array in which to store the FSR fission rates - * @param FSR_materials an array of FSR material UIDs + * @param FSR_materials an array of FSR material indices * @param materials an array of dev_material pointers * @param scalar_flux an array of FSR scalar fluxes */ @@ -720,11 +721,11 @@ GPUSolver::GPUSolver(Geometry* geometry, TrackGenerator* track_generator) : _scatter = NULL; _leakage = NULL; - if (track_generator != NULL) - setTrackGenerator(track_generator); - if (geometry != NULL) setGeometry(geometry); + + if (track_generator != NULL) + setTrackGenerator(track_generator); } @@ -1098,20 +1099,24 @@ void GPUSolver::initializeFSRs() { /* Allocate memory for all FSR volumes and dev_materials on the device */ try{ - /* Allocate memory on device for FSR volumes and Material UIDs */ + /* Allocate memory on device for FSR volumes and Material indices */ cudaMalloc((void**)&_FSR_volumes, _num_FSRs * sizeof(FP_PRECISION)); cudaMalloc((void**)&_FSR_materials, _num_FSRs * sizeof(int)); /* Create a temporary FSR array to populate and then copy to device */ FP_PRECISION* temp_FSR_volumes = new FP_PRECISION[_num_FSRs]; - /* Create a temporary FSR Material UIDs array to populate and then copy to device */ - int* FSRs_to_material_UIDs = new int[_num_FSRs]; + /* Create a temporary FSR to material indices array to populate and then + * copy to device */ + int* FSRs_to_material_indices = new int[_num_FSRs]; + + /* Initialize num fissionable FSRs counter */ _num_fissionable_FSRs = 0; - /* Populate FSR Material UIDs array */ + /* Populate FSR Material indices array */ for (int i = 0; i < _num_FSRs; i++){ - FSRs_to_material_UIDs[i] = _geometry->findFSRMaterial(i)->getUid(); + FSRs_to_material_indices[i] = _material_IDs_to_indices[_geometry-> + findFSRMaterial(i)->getId()]; if (_geometry->findFSRMaterial(i)->isFissionable()) _num_fissionable_FSRs++; } @@ -1148,18 +1153,18 @@ void GPUSolver::initializeFSRs() { /* Copy the temporary array of FSRs to the device */ cudaMemcpy((void*)_FSR_volumes, (void*)temp_FSR_volumes, _num_FSRs * sizeof(FP_PRECISION), cudaMemcpyHostToDevice); - cudaMemcpy((void*)_FSR_materials, (void*)FSRs_to_material_UIDs, + cudaMemcpy((void*)_FSR_materials, (void*)FSRs_to_material_indices, _num_FSRs * sizeof(int), cudaMemcpyHostToDevice); /* Copy the number of FSRs into constant memory on the GPU */ cudaMemcpyToSymbol(num_FSRs, (void*)&_num_FSRs, sizeof(int), 0, cudaMemcpyHostToDevice); - /* Free the temporary array of FSRs on the host */ + /* Free the temporary array of FSR volumes on the host */ free(temp_FSR_volumes); - /* Free the temporary array of FSR Material IDs on the host */ - free(FSRs_to_material_UIDs); + /* Free the temporary array of FSRs to material indices on the host */ + free(FSRs_to_material_indices); } catch(std::exception &e) { log_printf(ERROR, "Could not allocate memory for the GPUSolver's FSRs " @@ -1172,6 +1177,12 @@ void GPUSolver::initializeFSRs() { /** * @brief Allocates data on the GPU for all Materials data. + * @details This method loops over the materials in the host_materials map. + * Since cuda does not support std::map data types on the device (GPU), + * the materials map must be converted to an array and a map created + * that maps a material ID to an indice in the new materials array. In + * initializeTracks, this map is used to convert the Material ID + * associated with every segment to an index in the materials array. */ void GPUSolver::initializeMaterials() { @@ -1186,12 +1197,16 @@ void GPUSolver::initializeMaterials() { std::map host_materials=_geometry->getAllMaterials(); std::map::iterator iter; - + int material_index = 0; + /* Iterate through all Materials and clone them as dev_material structs * on the device */ cudaMalloc((void**)&_materials, _num_materials * sizeof(dev_material)); - for (iter=host_materials.begin(); iter != host_materials.end(); ++iter) - clone_material_on_gpu(iter->second, &_materials[iter->second->getUid()]); + for (iter=host_materials.begin(); iter != host_materials.end(); ++iter){ + clone_material_on_gpu(iter->second, &_materials[material_index]); + _material_IDs_to_indices[iter->second->getId()] = material_index; + material_index++; + } } catch(std::exception &e) { log_printf(ERROR, "Could not allocate memory for the GPUSolver's " @@ -1222,7 +1237,7 @@ void GPUSolver::initializeTracks() { for (int i=0; i < _tot_num_tracks; i++) { - clone_track_on_gpu(_tracks[i], &_dev_tracks[i]); + clone_track_on_gpu(_tracks[i], &_dev_tracks[i], _material_IDs_to_indices); /* Make Track reflective */ index = computeScalarTrackIndex(_tracks[i]->getTrackInI(), diff --git a/src/accel/cuda/GPUSolver.h b/src/accel/cuda/GPUSolver.h index b5f8d98a5..6cd479667 100644 --- a/src/accel/cuda/GPUSolver.h +++ b/src/accel/cuda/GPUSolver.h @@ -104,6 +104,9 @@ class GPUSolver : public Solver { /** Thrust vector of leakages for each track */ thrust::device_vector _leakage_vec; + /** Map of Material IDs to indices in _materials array */ + std::map _material_IDs_to_indices; + void initializePolarQuadrature(); void initializeFSRs(); void initializeMaterials(); diff --git a/src/accel/cuda/clone.cu b/src/accel/cuda/clone.cu index ec2ae8bbe..f9d1a2c68 100644 --- a/src/accel/cuda/clone.cu +++ b/src/accel/cuda/clone.cu @@ -12,15 +12,12 @@ */ void clone_material_on_gpu(Material* material_h, dev_material* material_d) { - /* Copy over the Material's ID and UID */ + /* Copy over the Material's ID */ int id = material_h->getId(); - int uid = material_h->getUid(); int num_groups = material_h->getNumEnergyGroups(); cudaMemcpy((void*)&material_d->_id, (void*)&id, sizeof(int), cudaMemcpyHostToDevice); - cudaMemcpy((void*)&material_d->_uid, (void*)&uid, sizeof(int), - cudaMemcpyHostToDevice); /* Allocate memory on the device for each dev_material data array */ double* sigma_t; @@ -71,15 +68,20 @@ void clone_material_on_gpu(Material* material_h, dev_material* material_d) { /** - * @brief Given a pointer to a Track on the host and a dev_track on - * the GPU, copy all of the class attributes and segments from - * the Track object on the host to the GPU. @details This - * routine is called by the GPUSolver::initializeTracks() - * private class method and is not intended to be called - * directly. @param track_h pointer to a Track on the host - * @param track_d pointer to a dev_track on the GPU + * @brief Given a pointer to a Track on the host, a dev_track on + * the GPU, and the map of material IDs to indices in the + * _materials array, copy all of the class attributes and + * segments from the Track object on the host to the GPU. + * @details This routine is called by the GPUSolver::initializeTracks() + * private class method and is not intended to be called + * directly. + * @param track_h pointer to a Track on the host + * @param track_d pointer to a dev_track on the GPU + * @param material_IDs_to_indices map of material IDs to indices + * in the _materials array. */ -void clone_track_on_gpu(Track* track_h, dev_track* track_d) { +void clone_track_on_gpu(Track* track_h, dev_track* track_d, + std::map &material_IDs_to_indices) { dev_segment* dev_segments; dev_segment* host_segments = new dev_segment[track_h->getNumSegments()]; @@ -101,7 +103,8 @@ void clone_track_on_gpu(Track* track_h, dev_track* track_d) { segment* curr = track_h->getSegment(s); host_segments[s]._length = curr->_length; host_segments[s]._region_uid = curr->_region_id; - host_segments[s]._material_uid = curr->_material->getUid(); + host_segments[s]._material_index = + material_IDs_to_indices[curr->_material->getId()]; } cudaMemcpy((void*)dev_segments, (void*)host_segments, diff --git a/src/accel/cuda/clone.h b/src/accel/cuda/clone.h index 1cab322ca..3f5a972b7 100644 --- a/src/accel/cuda/clone.h +++ b/src/accel/cuda/clone.h @@ -8,6 +8,8 @@ #include "../DeviceMaterial.h" #include "../DeviceTrack.h" +#include void clone_material_on_gpu(Material* material_h, dev_material* material_d); -void clone_track_on_gpu(Track* track_h, dev_track* track_d); +void clone_track_on_gpu(Track* track_h, dev_track* track_d, + std::map &material_IDs_to_indices);