Various tests and fixes of bugs exposed by them

Also fixed a high number of copy-paste bugs, so that the test passes.

Added a new component to the PR, to support N=0 element access to match a different current PR. I started this, before realising it will need some thought/rearchitecture to support.

Added HostAPI getVertex/EdgeCount(), and tested all current HostAPI functionality.

Expand tests to cover DeviceAPI.

Added support for not specifying length for host element methods.

Not sure I like this approach, it involved alot more internal changes than desired.

I was trying to avoid doing a double lookup, or returning a pair.

Enable (and fix) all the tests that were disabled due to not supporting no length passed to host element methods.

BugFix: DeviceCurve error string was not being passed right args.

Fixed how submodels bind curve for directed graphs.

This took a couple of attempts, I'm still not sure that it will support partial inheritance of a graph, that requires testing.

Initial (FULL MAPPING) submodel tests.

Still need to check how partial mapping would be handled, not sure we support that

Python Tests (and many bug fixes)

The Python HostException test is causing a hard to track down windows access violation.

Comitting so I can see if it segfaults on Linux.

Partial fix for that Python Access Violation

With GFlags, managed to find that the access violation was being triggered when Python is shutting down and the EnvironmentDirectedGraph::Host object is being destroyed late (because an exception was thrown and handled?).

So I switched a reference to a weak pointer, this solved 1 access violation but there's another in space between python and pyflamegpu with no debug symbols.

Parallel CSR construction.

Requires some thorough testing, I have manually tested that PBM construction operates as expected, however would be best to write a full test that checks edges are sorted correctly too

Rqures writing some edge traversal stuff too

Expect by sorting edges tests may now be broken

Fix docs CI

Can now iterate edges leaving a vertex.

Need to update/fix earlier tests now edges are being sorted.

Also need to update host copy of edges post-sort too

Syntax err, oddly only noticed by docs ci

This should fix the issue if host/device buffers not being sync'd

More tests now failing, due to the edge sort. Next job is to fix those.

include/flamegpu/defines.h

@@ -12,7 +12,15 @@ typedef unsigned int id_t;
  * Internal variable name used for IDs
  */
 constexpr const char* ID_VARIABLE_NAME = "_id";
+/**
+ * Internal variable name used for source-dest pairs
+ * @note These are always stored in [dest, source] order, but shown to the user in [source, dest] order. This enables 2D sorting
+ */
 constexpr const char* GRAPH_SOURCE_DEST_VARIABLE_NAME = "_source_dest";
+/**
+ * Internal variable name used to the pointer to the PBM of edges leaving each vertex
+ */
+constexpr const char* GRAPH_VERTEX_PBM_VARIABLE_NAME = "_pbm";
 /**
  * Internal value used when IDs have not be set
  * If this value is changed, things may break

include/flamegpu/gpu/CUDAAgent.h

@@ -42,7 +42,7 @@ class CUDAAgent : public AgentInterface {
      *  map of agent function name to RTC function instance
      */
      typedef std::map<const std::string, std::unique_ptr<jitify::experimental::KernelInstantiation>> CUDARTCFuncMap;
-     typedef std::map<const std::string, std::unique_ptr<detail::curve::CurveRTCHost>> CUDARTCHeaderMap;
+     typedef std::map<const std::string, std::shared_ptr<detail::curve::CurveRTCHost>> CUDARTCHeaderMap;
     /**
      * Element type of CUDARTCFuncMap
      */
@@ -384,7 +384,7 @@ class CUDAAgent : public AgentInterface {
      * map between function name (or function_name_condition) and the curve interface
      * This allows access for updating curve
      */
-    std::unordered_map<std::string, std::unique_ptr<detail::curve::HostCurve>> curve_map;
+    std::unordered_map<std::string, std::shared_ptr<detail::curve::HostCurve>> curve_map;
     /**
      * Used when allocated new buffers
      */

include/flamegpu/gpu/CUDAScatter.cuh

@@ -147,6 +147,22 @@ class CUDAScatter {
         Type messageOrAgent,
         const std::vector<ScatterData>& scatterData,
         unsigned int itemCount);
+    /**
+     * Scatters agents from SoA to SoA according to d_position flag as input_source, all variables are scattered
+     * Used for Host function sort agent
+     * CUDAScanCompaction::position is used to decide where to scatter to
+     * @param streamResourceId The stream index to use for accessing stream specific resources such as scan compaction arrays and buffers
+     * @param stream CUDA stream to be used for async CUDA operations
+     * @param position Buffer containing indexes to move data from
+     * @param scatterData Vector of scatter configuration for each variable to be scattered
+     * @param itemCount Total number of items in input array to consider
+     */
+    void scatterPosition_async(
+        unsigned int streamResourceId,
+        cudaStream_t stream,
+        unsigned int *position,
+        const std::vector<ScatterData>& scatterData,
+        unsigned int itemCount);
     void scatterPosition(
         unsigned int streamResourceId,
         cudaStream_t stream,

include/flamegpu/gpu/CUDASimulation.h

@@ -17,6 +17,7 @@
 #include "flamegpu/runtime/HostNewAgentAPI.h"
 #include "flamegpu/gpu/CUDAMacroEnvironment.h"
 #include "flamegpu/runtime/utility/EnvironmentManager.cuh"
+#include "flamegpu/runtime/environment/EnvironmentDirectedGraph.cuh"
 
 #ifdef FLAMEGPU_VISUALISATION
 #include "flamegpu/visualiser/ModelVis.h"

include/flamegpu/model/EnvironmentDescription.h

@@ -13,6 +13,7 @@
 #include "flamegpu/model/EnvironmentData.h"
 #include "flamegpu/util/type_decode.h"
 #include "flamegpu/gpu/CUDAEnsemble.h"
+#include "flamegpu/runtime/environment/EnvironmentDirectedGraph.cuh"
 #include "flamegpu/runtime/environment/EnvironmentDirectedGraph/EnvironmentDirectedGraphDescription.cuh"
 
 namespace flamegpu {

include/flamegpu/runtime/detail/curve/DeviceCurve.cuh

@@ -232,6 +232,15 @@ class DeviceCurve {
     template <typename T, unsigned int N = 0, unsigned int M>
     __device__ __forceinline__ static T getEnvironmentArrayProperty(const char(&propertyName)[M], unsigned int array_index);
 
+
+    /**
+     * Retrieve the address of the specified environment directed graph's pbm for edges leaving each vertex
+     *
+     * @param graphHash The hash of the affected graph
+     * @return The requested address
+     * @throws exception::DeviceError (Only when SEATBELTS==ON) If the specified graph is not found in the cuRVE hashtable, or it's details are invalid.
+     */
+    __device__ __forceinline__ static unsigned int *getEnvironmentDirectedGraphPBM(VariableHash graphHash);
     /**
      * Retrieve the specified environment directed graph's vertex (array) property from the cuRVE hashtable.
      *
@@ -315,7 +324,7 @@ __device__ __forceinline__ char* DeviceCurve::getVariablePtr(const char(&variabl
         DTHROW("Curve variable with name '%s', variable array length mismatch %u != %u.\n", variableName, sm()->curve_elements[cv], type_decode<T>::len_t);
         return nullptr;
     } else if (offset >= sm()->curve_type_size[cv] * sm()->curve_elements[cv] * sm()->curve_count[cv]) {  // Note : offset is basically index * sizeof(T)
-        DTHROW("Curve variable with name '%s', offset exceeds buffer length  %u >= %u.\n", offset, sm()->curve_type_size[cv] * sm()->curve_elements[cv] * sm()->curve_count[cv]);
+        DTHROW("Curve variable with name '%s', offset exceeds buffer length  %u >= %u.\n", variableName, offset, sm()->curve_type_size[cv] * sm()->curve_elements[cv] * sm()->curve_count[cv]);
         return nullptr;
     }
 #endif
@@ -433,6 +442,10 @@ __device__ __forceinline__ T DeviceCurve::getEnvironmentArrayProperty(const char
     using detail::sm;
     return *reinterpret_cast<const T*>(sm()->env_buffer + reinterpret_cast<ptrdiff_t>(getVariablePtr<T, N>(propertyName, Curve::variableHash("_environment"), array_index * sizeof(T))));
 }
+__device__ __forceinline__ unsigned int *DeviceCurve::getEnvironmentDirectedGraphPBM(VariableHash graphHash) {
+    using detail::sm;
+    return static_cast<unsigned int *>(static_cast<void*>(getVariablePtr<unsigned int, 1>("_pbm", graphHash ^ Curve::variableHash("_environment_directed_graph_vertex"), 0)));
+}
 template <typename T, unsigned int M>
 __device__ __forceinline__ T DeviceCurve::getEnvironmentDirectedGraphVertexProperty(VariableHash graphHash, const char(&propertyName)[M], unsigned int vertex_index) {
     using detail::sm;

include/flamegpu/runtime/environment/EnvironmentDirectedGraph.cuh

@@ -1,8 +1,6 @@
 #ifndef INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_ENVIRONMENTDIRECTEDGRAPH_CUH_
 #define INCLUDE_FLAMEGPU_RUNTIME_ENVIRONMENT_ENVIRONMENTDIRECTEDGRAPH_CUH_
 
-#include "flamegpu/runtime/utility/EnvironmentManager.cuh"
-
 namespace flamegpu {
 /**
  * Environment Directed Graph functionality