diff --git a/CHANGELOG.md b/CHANGELOG.md
index d55ce40ac5b3..978bccbf19c0 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -3,6 +3,7 @@
 ## Current develop
 
 ### Added (new features/APIs/variables/...)
+- [[PR 907]](https://github.com/parthenon-hpc-lab/parthenon/pull/907) PEP1: Allow subclassing StateDescriptor
 - [[PR 932]](https://github.com/parthenon-hpc-lab/parthenon/pull/932) Add GetOrAddFlag to metadata
 - [[PR 931]](https://github.com/parthenon-hpc-lab/parthenon/pull/931) Allow SparsePacks with subsets of blocks
 - [[PR 921]](https://github.com/parthenon-hpc-lab/parthenon/pull/921) Add more flexible ways of adding and using MeshData/MeshBlockData objects to DataCollections
@@ -35,6 +36,7 @@
 - [[PR 890]](https://github.com/parthenon-hpc-lab/parthenon/pull/890) Fix bugs in sparse communication and prolongation
 
 ### Infrastructure (changes irrelevant to downstream codes)
+- [[PR 967]](https://github.com/parthenon-hpc-lab/parthenon/pull/967) Change INLINE to FORCEINLINE on par_for_inner overloads
 - [[PR 938]](https://github.com/parthenon-hpc-lab/parthenon/pull/938) Restructure buffer packing/unpacking kernel hierarchical parallelism
 - [[PR 944]](https://github.com/parthenon-hpc-lab/parthenon/pull/944) Move sparse pack identifier creation to descriptor
 - [[PR 904]](https://github.com/parthenon-hpc-lab/parthenon/pull/904) Move to prolongation/restriction in one for AMR and communicate non-cell centered fields
diff --git a/src/interface/packages.hpp b/src/interface/packages.hpp
index 285a815bf9ea..1ef6db59f58f 100644
--- a/src/interface/packages.hpp
+++ b/src/interface/packages.hpp
@@ -16,6 +16,7 @@
 
 #include <memory>
 #include <string>
+#include <vector>
 
 #include "basic_types.hpp"
 
@@ -26,15 +27,47 @@ class Packages_t {
   Packages_t() = default;
   void Add(const std::shared_ptr<StateDescriptor> &package);
 
-  std::shared_ptr<StateDescriptor> const &Get(const std::string &name) {
+  std::shared_ptr<StateDescriptor> const &Get(const std::string &name) const {
     return packages_.at(name);
   }
 
+  // Retrieve a package pointer, cast to a given type T
+  template <typename T>
+  T *Get(const std::string &name) const {
+    return static_cast<T *>(packages_.at(name).get());
+  }
+
   const Dictionary<std::shared_ptr<StateDescriptor>> &AllPackages() const {
     return packages_;
   }
   Dictionary<std::shared_ptr<StateDescriptor>> &AllPackages() { return packages_; }
 
+  // Returns a sub-Dictionary containing just pointers to packages of type T.
+  // Dictionary is a *new copy*, and members are bare pointers, not shared_ptr.
+  template <typename T>
+  const Dictionary<T *> AllPackagesOfType() const {
+    Dictionary<T *> sub_dict;
+    for (auto package : packages_) {
+      if (T *cast_package = dynamic_cast<T *>(package.second.get())) {
+        sub_dict[package.first] = cast_package;
+      }
+    }
+    return sub_dict;
+  }
+
+  // Returns a list of pointers to packages of type T.
+  // List contains bare pointers, not shared_ptr objects
+  template <typename T>
+  const std::vector<T *> ListPackagesOfType() const {
+    std::vector<T *> sub_list;
+    for (auto package : packages_) {
+      if (T *cast_package = dynamic_cast<T *>(package.second.get())) {
+        sub_list.append(cast_package);
+      }
+    }
+    return sub_list;
+  }
+
  private:
   Dictionary<std::shared_ptr<StateDescriptor>> packages_;
 };
diff --git a/src/interface/state_descriptor.hpp b/src/interface/state_descriptor.hpp
index ff21e628147c..5bb37ea83b60 100644
--- a/src/interface/state_descriptor.hpp
+++ b/src/interface/state_descriptor.hpp
@@ -108,6 +108,9 @@ class StateDescriptor {
     }
   }
 
+  // Virtual destructor for subclassing
+  virtual ~StateDescriptor() = default;
+
   static std::shared_ptr<StateDescriptor>
   CreateResolvedStateDescriptor(Packages_t &packages);
 
@@ -429,7 +432,7 @@ class StateDescriptor {
 
   friend std::ostream &operator<<(std::ostream &os, const StateDescriptor &sd);
 
- private:
+ protected:
   void InvertControllerMap();
 
   Params params_;
diff --git a/src/kokkos_abstraction.hpp b/src/kokkos_abstraction.hpp
index cabae1273b6f..bfaf6a247a66 100644
--- a/src/kokkos_abstraction.hpp
+++ b/src/kokkos_abstraction.hpp
@@ -1,9 +1,9 @@
 //========================================================================================
 // Parthenon performance portable AMR framework
-// Copyright(C) 2020-2022 The Parthenon collaboration
+// Copyright(C) 2020-2023 The Parthenon collaboration
 // Licensed under the 3-clause BSD License, see LICENSE file for details
 //========================================================================================
-// (C) (or copyright) 2020-2022. Triad National Security, LLC. All rights reserved.
+// (C) (or copyright) 2020-2023. Triad National Security, LLC. All rights reserved.
 //
 // This program was produced under U.S. Government contract 89233218CNA000001
 // for Los Alamos National Laboratory (LANL), which is operated by Triad
@@ -707,7 +707,7 @@ inline void par_for_outer(OuterLoopPatternTeams, const std::string &name,
 
 // Inner parallel loop using TeamThreadRange
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void
+KOKKOS_FORCEINLINE_FUNCTION void
 par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ll, const int lu,
               const int ml, const int mu, const int nl, const int nu, const int kl,
               const int ku, const int jl, const int ju, const int il, const int iu,
@@ -741,7 +741,7 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ll, const i
       });
 }
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void
+KOKKOS_FORCEINLINE_FUNCTION void
 par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ml, const int mu,
               const int nl, const int nu, const int kl, const int ku, const int jl,
               const int ju, const int il, const int iu, const Function &function) {
@@ -770,7 +770,7 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int ml, const i
                        });
 }
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void
+KOKKOS_FORCEINLINE_FUNCTION void
 par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int nl, const int nu,
               const int kl, const int ku, const int jl, const int ju, const int il,
               const int iu, const Function &function) {
@@ -795,10 +795,10 @@ par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int nl, const i
                        });
 }
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member,
-                                          const int kl, const int ku, const int jl,
-                                          const int ju, const int il, const int iu,
-                                          const Function &function) {
+KOKKOS_FORCEINLINE_FUNCTION void
+par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int kl, const int ku,
+              const int jl, const int ju, const int il, const int iu,
+              const Function &function) {
   const int Nk = ku - kl + 1;
   const int Nj = ju - jl + 1;
   const int Ni = iu - il + 1;
@@ -815,9 +815,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_m
   });
 }
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member,
-                                          const int jl, const int ju, const int il,
-                                          const int iu, const Function &function) {
+KOKKOS_FORCEINLINE_FUNCTION void
+par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member, const int jl, const int ju,
+              const int il, const int iu, const Function &function) {
   const int Nj = ju - jl + 1;
   const int Ni = iu - il + 1;
   const int NjNi = Nj * Ni;
@@ -828,22 +828,22 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_m
   });
 }
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR, team_mbr_t team_member,
-                                          const int il, const int iu,
-                                          const Function &function) {
+KOKKOS_FORCEINLINE_FUNCTION void par_for_inner(InnerLoopPatternTTR,
+                                               team_mbr_t team_member, const int il,
+                                               const int iu, const Function &function) {
   Kokkos::parallel_for(Kokkos::TeamThreadRange(team_member, il, iu + 1), function);
 }
 // Inner parallel loop using TeamVectorRange
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternTVR, team_mbr_t team_member,
-                                          const int il, const int iu,
-                                          const Function &function) {
+KOKKOS_FORCEINLINE_FUNCTION void par_for_inner(InnerLoopPatternTVR,
+                                               team_mbr_t team_member, const int il,
+                                               const int iu, const Function &function) {
   Kokkos::parallel_for(Kokkos::TeamVectorRange(team_member, il, iu + 1), function);
 }
 
 // Inner parallel loop using FOR SIMD
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void
+KOKKOS_FORCEINLINE_FUNCTION void
 par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int nl, const int nu,
               const int kl, const int ku, const int jl, const int ju, const int il,
               const int iu, const Function &function) {
@@ -859,10 +859,10 @@ par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int nl, con
   }
 }
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member,
-                                          const int kl, const int ku, const int jl,
-                                          const int ju, const int il, const int iu,
-                                          const Function &function) {
+KOKKOS_FORCEINLINE_FUNCTION void
+par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int kl, const int ku,
+              const int jl, const int ju, const int il, const int iu,
+              const Function &function) {
   for (int k = kl; k <= ku; ++k) {
     for (int j = jl; j <= ju; ++j) {
 #pragma omp simd
@@ -873,9 +873,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t te
   }
 }
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member,
-                                          const int jl, const int ju, const int il,
-                                          const int iu, const Function &function) {
+KOKKOS_FORCEINLINE_FUNCTION void
+par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member, const int jl, const int ju,
+              const int il, const int iu, const Function &function) {
   for (int j = jl; j <= ju; ++j) {
 #pragma omp simd
     for (int i = il; i <= iu; i++) {
@@ -884,9 +884,9 @@ KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t te
   }
 }
 template <typename Function>
-KOKKOS_INLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor, team_mbr_t team_member,
-                                          const int il, const int iu,
-                                          const Function &function) {
+KOKKOS_FORCEINLINE_FUNCTION void par_for_inner(InnerLoopPatternSimdFor,
+                                               team_mbr_t team_member, const int il,
+                                               const int iu, const Function &function) {
 #pragma omp simd
   for (int i = il; i <= iu; i++) {
     function(i);