diff --git a/include/openmc/message_passing.h b/include/openmc/message_passing.h
index c730674ab4..82b27a7790 100644
--- a/include/openmc/message_passing.h
+++ b/include/openmc/message_passing.h
@@ -13,7 +13,7 @@ namespace mpi {
   extern bool master;
 
 #ifdef OPENMC_MPI
-  extern MPI_Datatype bank;
+  extern MPI_Datatype source_site;
   extern MPI_Comm intracomm;
 #endif
 
diff --git a/include/openmc/particle.h b/include/openmc/particle.h
index 2d38bd2732..0fda3d3079 100644
--- a/include/openmc/particle.h
+++ b/include/openmc/particle.h
@@ -21,11 +21,12 @@ namespace openmc {
 // Forward declare the Surface class for use in Particle::cross_vacuum_bc, etc.
 class Surface;
 
-//============================================================================
-//! State of a particle being transported through geometry
-//! This class defines actions particles can take. Its base
-//! class defines particle data layout in memory.
-//============================================================================
+/*
+ * The Particle class encompasses data and methods for transporting particles
+ * through their lifecycle. Its base class defines particle data layout in
+ * memory. A more detailed description of the rationale behind this approach
+ * can be found in particle_data.h.
+ */
 
 class Particle : public ParticleData {
 public:
diff --git a/include/openmc/particle_data.h b/include/openmc/particle_data.h
index 6050ca96e1..d3d745b4ed 100644
--- a/include/openmc/particle_data.h
+++ b/include/openmc/particle_data.h
@@ -164,6 +164,32 @@ struct BoundaryInfo {
 //! Defines how particle data is laid out in memory
 //============================================================================
 
+/*
+ * This class was added in order to separate the layout and access of particle
+ * data from particle physics operations during a development effort to get
+ * OpenMC running on GPUs. In the event-based Monte Carlo method, one creates
+ * an array of particles on which actions like cross section lookup and surface
+ * crossing are done en masse, which works best on vector computers of yore and
+ * modern GPUs. It has been shown in the below publication that arranging
+ * particle data into a structure of arrays rather than an array of structures
+ * enhances performance on GPUs. For instance, rather than having an
+ * std::vector<Particle> where consecutive particle energies would be separated
+ * by about 400 bytes, one would create a structure which has a single
+ * std::vector<double> of energies.  The motivation here is that more coalesced
+ * memory accesses occur, in the parlance of GPU programming.
+ *
+ * So, this class enables switching between the array-of-structures and
+ * structure- of-array data layout at compile time. In GPU branches of the
+ * code, our Particle class inherits from a class that provides an array of
+ * particle energies, and can access them using the E() method (defined below).
+ * In the CPU code, we inherit from this class which gives the conventional
+ * layout of particle data, useful for history-based tracking.
+ *
+ * As a result, we always use the E(), r_last(), etc. methods to access
+ * particle data in order to keep a unified interface between
+ * structure-of-array and array-of-structure code on either CPU or GPU code
+ * while sharing the same physics code on each codebase.
+ */
 class ParticleData {
 
 public:
diff --git a/openmc/lib/core.py b/openmc/lib/core.py
index 6d9a680e35..e740e28699 100644
--- a/openmc/lib/core.py
+++ b/openmc/lib/core.py
@@ -11,7 +11,7 @@ from .error import _error_handler
 import openmc.lib
 
 
-class _Bank(Structure):
+class _SourceSite(Structure):
     _fields_ = [('r', c_double*3),
                 ('u', c_double*3),
                 ('E', c_double),
@@ -73,7 +73,7 @@ _run_linsolver_argtypes = [_array_1d_dble, _array_1d_dble, _array_1d_dble,
                            c_double]
 _dll.openmc_run_linsolver.argtypes = _run_linsolver_argtypes
 _dll.openmc_run_linsolver.restype = c_int
-_dll.openmc_source_bank.argtypes = [POINTER(POINTER(_Bank)), POINTER(c_int64)]
+_dll.openmc_source_bank.argtypes = [POINTER(POINTER(_SourceSite)), POINTER(c_int64)]
 _dll.openmc_source_bank.restype = c_int
 _dll.openmc_source_bank.errcheck = _error_handler
 _dll.openmc_simulation_init.restype = c_int
@@ -342,13 +342,13 @@ def source_bank():
 
     """
     # Get pointer to source bank
-    ptr = POINTER(_Bank)()
+    ptr = POINTER(_SourceSite)()
     n = c_int64()
     _dll.openmc_source_bank(ptr, n)
 
     try:
         # Convert to numpy array with appropriate datatype
-        bank_dtype = np.dtype(_Bank)
+        bank_dtype = np.dtype(_SourceSite)
         return as_array(ptr, (n.value,)).view(bank_dtype)
 
     except ValueError as err:
diff --git a/src/eigenvalue.cpp b/src/eigenvalue.cpp
index 84d5fec0bb..1e0d7d972e 100644
--- a/src/eigenvalue.cpp
+++ b/src/eigenvalue.cpp
@@ -230,7 +230,7 @@ void synchronize_bank()
       // process
       if (neighbor != mpi::rank) {
         requests.emplace_back();
-        MPI_Isend(&temp_sites[index_local], static_cast<int>(n), mpi::bank,
+        MPI_Isend(&temp_sites[index_local], static_cast<int>(n), mpi::source_site,
           neighbor, mpi::rank, mpi::intracomm, &requests.back());
       }
 
@@ -276,7 +276,7 @@ void synchronize_bank()
       // asynchronous receive for the source sites
 
       requests.emplace_back();
-      MPI_Irecv(&simulation::source_bank[index_local], static_cast<int>(n), mpi::bank,
+      MPI_Irecv(&simulation::source_bank[index_local], static_cast<int>(n), mpi::source_site,
             neighbor, neighbor, mpi::intracomm, &requests.back());
 
     } else {
diff --git a/src/finalize.cpp b/src/finalize.cpp
index 45dd0fc7cc..22dc671836 100644
--- a/src/finalize.cpp
+++ b/src/finalize.cpp
@@ -143,7 +143,7 @@ int openmc_finalize()
 
   // Free all MPI types
 #ifdef OPENMC_MPI
-  if (mpi::bank != MPI_DATATYPE_NULL) MPI_Type_free(&mpi::bank);
+  if (mpi::source_site != MPI_DATATYPE_NULL) MPI_Type_free(&mpi::source_site);
 #endif
 
   return 0;
diff --git a/src/initialize.cpp b/src/initialize.cpp
index c6a432d125..d67d6b3b31 100644
--- a/src/initialize.cpp
+++ b/src/initialize.cpp
@@ -151,8 +151,8 @@ void initialize_mpi(MPI_Comm intracomm)
 
   int blocks[] {3, 3, 1, 1, 1, 1, 1, 1, 1};
   MPI_Datatype types[] {MPI_DOUBLE, MPI_DOUBLE, MPI_DOUBLE, MPI_DOUBLE, MPI_INT, MPI_INT, MPI_INT, MPI_LONG, MPI_LONG};
-  MPI_Type_create_struct(9, blocks, disp, types, &mpi::bank);
-  MPI_Type_commit(&mpi::bank);
+  MPI_Type_create_struct(9, blocks, disp, types, &mpi::source_site);
+  MPI_Type_commit(&mpi::source_site);
 }
 #endif // OPENMC_MPI
 
diff --git a/src/message_passing.cpp b/src/message_passing.cpp
index aea17f9f49..f479d3db74 100644
--- a/src/message_passing.cpp
+++ b/src/message_passing.cpp
@@ -9,7 +9,7 @@ bool master {true};
 
 #ifdef OPENMC_MPI
 MPI_Comm intracomm {MPI_COMM_NULL};
-MPI_Datatype bank {MPI_DATATYPE_NULL};
+MPI_Datatype source_site {MPI_DATATYPE_NULL};
 #endif
 
 extern "C" bool openmc_master() { return mpi::master; }
diff --git a/src/state_point.cpp b/src/state_point.cpp
index 61c2a5339e..f98727a08d 100644
--- a/src/state_point.cpp
+++ b/src/state_point.cpp
@@ -666,7 +666,7 @@ write_source_bank(hid_t group_id, bool surf_source_bank)
 #ifdef OPENMC_MPI
       // Receive source sites from other processes
       if (i > 0)
-        MPI_Recv(source_bank->data(), count[0], mpi::bank, i, i,
+        MPI_Recv(source_bank->data(), count[0], mpi::source_site, i, i,
                  mpi::intracomm, MPI_STATUS_IGNORE);
 #endif
 
@@ -691,7 +691,7 @@ write_source_bank(hid_t group_id, bool surf_source_bank)
 #endif
   } else {
 #ifdef OPENMC_MPI
-    MPI_Send(source_bank->data(), count_size, mpi::bank,
+    MPI_Send(source_bank->data(), count_size, mpi::source_site,
       0, mpi::rank, mpi::intracomm);
 #endif
   }