Most changes to get EnvironmentManager refactor working.

3 tests inside SubEnvironmentManagerTest failing. Not checked python tests yet, some may have broken due to API change.

Author: Robadob

include/flamegpu/exception/FLAMEGPUDeviceException_device.cuh

@@ -6,6 +6,8 @@
 
 #include <cstring>
 
+#include "flamegpu/runtime/detail/SharedBlock.h"
+
 namespace flamegpu {
 namespace exception {
 
@@ -81,20 +83,20 @@ class DeviceException {
      */
     template<typename... Args>
     __device__ void setMessage(const char *format, Args... args) {
-        extern __shared__ DeviceExceptionBuffer* buff[];
+        using detail::sm;
         // Only the thread which first reported error gets to output
         if (hasError) {
             // Only output once
-            if (buff[2]->format_string[0])
+            if (sm()->device_exception->format_string[0])
                 return;
             // Copy the format string
             unsigned int eos = 0;
             for (eos = 0; eos < DeviceExceptionBuffer::FORMAT_BUFF_LEN; ++eos)
                 if (format[eos] == '\0')
                     break;
-            memcpy(buff[2]->format_string, format, eos * sizeof(char));
+            memcpy(sm()->device_exception->format_string, format, eos * sizeof(char));
             // Process args
-            subformat_recurse(buff[2], args...);
+            subformat_recurse(sm()->device_exception, args...);
         }
     }
 
@@ -135,18 +137,18 @@ class DeviceException {
      */
     __device__ DeviceException(const char *file, const unsigned int line)
         : hasError(!getErrorCount()) {
-        extern __shared__ DeviceExceptionBuffer* buff[];
+        using detail::sm;
         if (hasError) {
             // Copy file location
             const size_t file_len = strlen(file);
-            memcpy(buff[2]->file_path, file, file_len);
+            memcpy(sm()->device_exception->file_path, file, file_len);
             // Copy line no
-            buff[2]->line_no = line;
+            sm()->device_exception->line_no = line;
             // Copy block/thread indices
             const uint3 bid3 = blockIdx;
-            memcpy(buff[2]->block_id, &bid3, sizeof(unsigned int) * 3);
+            memcpy(sm()->device_exception->block_id, &bid3, sizeof(unsigned int) * 3);
             const uint3 tid3 = threadIdx;
-            memcpy(buff[2]->thread_id, &tid3, sizeof(unsigned int) * 3);
+            memcpy(sm()->device_exception->thread_id, &tid3, sizeof(unsigned int) * 3);
         }
     }
     /**
@@ -188,9 +190,9 @@ __device__ inline void DeviceException::subformat(DeviceExceptionBuffer *buff, c
     }
 }
 __device__ unsigned int DeviceException::getErrorCount() {
-    extern __shared__ DeviceExceptionBuffer* buff[];
+    using detail::sm;
     // Are we the first exception
-    return atomicInc(&buff[2]->error_count, UINT_MAX);
+    return atomicInc(&sm()->device_exception->error_count, UINT_MAX);
 }
 #endif
 #else

include/flamegpu/io/JSONStateReader.h

@@ -33,7 +33,7 @@ class JSONStateReader : public StateReader {
     JSONStateReader(
         const std::string &model_name,
         const std::unordered_map<std::string, EnvironmentDescription::PropData> &env_desc,
-        util::StringUint32PairUnorderedMap<util::Any> &env_init,
+        std::unordered_map<std::string, util::Any>&env_init,
         util::StringPairUnorderedMap<std::shared_ptr<AgentVector>> &model_state,
         const std::string &input_file,
         Simulation *sim_instance);

include/flamegpu/io/StateReader.h

@@ -37,7 +37,7 @@ class StateReader {
     StateReader(
         const std::string& _model_name,
         const std::unordered_map<std::string, EnvironmentDescription::PropData>& _env_desc,
-        util::StringUint32PairUnorderedMap<util::Any>& _env_init,
+        std::unordered_map<std::string, util::Any>& _env_init,
         util::StringPairUnorderedMap<std::shared_ptr<AgentVector>>& _model_state,
         const std::string& input,
         Simulation* _sim_instance)
@@ -67,7 +67,7 @@ class StateReader {
     std::string inputFile;
     const std::string model_name;
     const std::unordered_map<std::string, EnvironmentDescription::PropData> &env_desc;
-    util::StringUint32PairUnorderedMap<util::Any> &env_init;
+    std::unordered_map<std::string, util::Any>& env_init;
     Simulation *sim_instance;
 };
 }  // namespace io

include/flamegpu/io/StateReaderFactory.h

@@ -39,7 +39,7 @@ class StateReaderFactory {
     static StateReader* createReader(
         const std::string& model_name,
         const std::unordered_map<std::string, EnvironmentDescription::PropData>& env_desc,
-        util::StringUint32PairUnorderedMap<util::Any>& env_init,
+        std::unordered_map<std::string, util::Any>& env_init,
         util::StringPairUnorderedMap<std::shared_ptr<AgentVector>>& model_state,
         const std::string& input,
         Simulation* sim_instance) {

include/flamegpu/io/XMLStateReader.h

@@ -32,7 +32,7 @@ class XMLStateReader : public StateReader {
     XMLStateReader(
         const std::string &model_name,
         const std::unordered_map<std::string, EnvironmentDescription::PropData> &env_desc,
-        util::StringUint32PairUnorderedMap<util::Any> &env_init,
+        std::unordered_map<std::string, util::Any> &env_init,
         util::StringPairUnorderedMap<std::shared_ptr<AgentVector>> &model_state,
         const std::string &input_file,
         Simulation *sim_instance);

include/flamegpu/runtime/AgentFunction.cuh

@@ -4,6 +4,7 @@
 #include <cuda_runtime.h>
 #include <curand_kernel.h>
 
+#include "flamegpu/runtime/detail/SharedBlock.h"
 #include "flamegpu/defines.h"
 #include "flamegpu/exception/FLAMEGPUDeviceException.cuh"
 #include "flamegpu/runtime/AgentFunction_shim.cuh"
@@ -18,7 +19,7 @@ typedef void(AgentFunctionWrapper)(
     exception::DeviceExceptionBuffer *error_buffer,
 #endif
 #ifndef __CUDACC_RTC__
-    const detail::curve::Curve::CurveTable *d_curve_table,
+    const detail::curve::CurveTable *d_curve_table,
     const char* d_env_buffer,
 #endif
     id_t *d_agent_output_nextID,
@@ -54,7 +55,7 @@ __global__ void agent_function_wrapper(
     exception::DeviceExceptionBuffer *error_buffer,
 #endif
 #ifndef __CUDACC_RTC__
-    const detail::curve::Curve::CurveTable* d_curve_table,
+    const detail::curve::CurveTable* d_curve_table,
     const char* d_env_buffer,
 #endif
     id_t *d_agent_output_nextID,
@@ -67,13 +68,13 @@ __global__ void agent_function_wrapper(
     unsigned int *scanFlag_agentOutput) {
     // We place these at the start of shared memory, so we can locate it anywhere in device code without a reference
     if (threadIdx.x == 0) {
-        extern __shared__ const void* sm[];
+        using detail::sm;
 #ifndef __CUDACC_RTC__
-        sm[0] = d_curve_table;
-        sm[1] = d_env_buffer;
+        sm()->curve = d_curve_table;
+        sm()->env_buffer = d_env_buffer;
 #endif
 #if !defined(SEATBELTS) || SEATBELTS
-        sm[2] = error_buffer;
+        sm()->device_exception = error_buffer;
 #endif
     }
 

include/flamegpu/runtime/AgentFunctionCondition.cuh

@@ -14,7 +14,7 @@ typedef void(AgentFunctionConditionWrapper)(
     exception::DeviceExceptionBuffer *error_buffer,
 #endif
 #ifndef __CUDACC_RTC__
-    const detail::curve::Curve::CurveTable* d_curve_table,
+    const detail::curve::CurveTable* d_curve_table,
     const char* d_env_buffer,
 #endif
     const unsigned int popNo,
@@ -39,21 +39,21 @@ __global__ void agent_function_condition_wrapper(
     exception::DeviceExceptionBuffer *error_buffer,
 #endif
 #ifndef __CUDACC_RTC__
-    const detail::curve::Curve::CurveTable* d_curve_table,
+    const detail::curve::CurveTable* d_curve_table,
     const char* d_env_buffer,
 #endif
     const unsigned int popNo,
     curandState *d_rng,
     unsigned int *scanFlag_conditionResult) {
     // We place these at the start of shared memory, so we can locate it anywhere in device code without a reference
     if (threadIdx.x == 0) {
-        extern __shared__ const void* sm[];
+        using detail::sm;
 #ifndef __CUDACC_RTC__
-        sm[0] = d_curve_table;
-        sm[1] = d_env_buffer;
+        sm()->curve = d_curve_table;
+        sm()->env_buffer = d_env_buffer;
 #endif
 #if !defined(SEATBELTS) || SEATBELTS
-        sm[2] = error_buffer;
+        sm()->device_exception = error_buffer;
 #endif
     }
 

include/flamegpu/runtime/DeviceAPI.cuh

@@ -44,7 +44,7 @@ class ReadOnlyDeviceAPI {
         exception::DeviceExceptionBuffer *,
 #endif
 #ifndef __CUDACC_RTC__
-        const detail::curve::Curve::CurveTable *,
+        const detail::curve::CurveTable *,
 #endif
         const unsigned int,
         curandState *,
@@ -148,7 +148,7 @@ class DeviceAPI {
         exception::DeviceExceptionBuffer *,
 #endif
 #ifndef __CUDACC_RTC__
-        const detail::curve::Curve::CurveTable *,
+        const detail::curve::CurveTable *,
 #endif
         id_t*,
         const unsigned int,
@@ -387,6 +387,7 @@ __device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[M], const
 template<typename MessageIn, typename MessageOut>
 template<typename T, unsigned int N>
 __device__ T DeviceAPI<MessageIn, MessageOut>::getVariable(const char(&variable_name)[N]) const {
+    using detail::sm;
     // simple indexing assumes index is the thread number (this may change later)
     const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x;
 

include/flamegpu/runtime/detail/SharedBlock.h

@@ -0,0 +1,35 @@
+#ifndef INCLUDE_FLAMEGPU_RUNTIME_DETAIL_SHAREDBLOCK_H_
+#define INCLUDE_FLAMEGPU_RUNTIME_DETAIL_SHAREDBLOCK_H_
+
+namespace flamegpu {
+namespace exception {
+struct DeviceExceptionBuffer;
+}  // namespace exception
+namespace detail {
+namespace curve {
+struct CurveTable;
+}  // namespace curve
+/**
+ * This struct represents the data we package into shared memory
+ * The ifndef __CUDACC_RTC__ will cause the size to be too large for RTC builds, but that's not (currently) an issue
+ */
+struct SharedBlock {
+#ifndef __CUDACC_RTC__
+    const curve::CurveTable* curve;
+    const char* env_buffer;
+#endif
+#if !defined(SEATBELTS) || SEATBELTS
+    exception::DeviceExceptionBuffer *device_exception;
+#endif
+};
+/**
+ * Returns a pointer to a common item in shared memory
+ */
+__forceinline__ __device__ SharedBlock *sm() {
+    __shared__ SharedBlock _sm;
+    return &_sm;
+}
+}  // namespace detail
+}  // namespace flamegpu
+
+#endif  // INCLUDE_FLAMEGPU_RUNTIME_DETAIL_SHAREDBLOCK_H_

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

@@ -29,13 +29,6 @@ class Curve {
     typedef unsigned int             NamespaceHash;      // !< Typedef for cuRVE variable namespace string hash
     static const int MAX_VARIABLES = 1024;               // !< Default maximum number of cuRVE variables (must be a power of 2)
     static const VariableHash EMPTY_FLAG = 0;
-    struct CurveTable {
-        VariableHash hashes[MAX_VARIABLES];            // Device array of the hash values of registered variables
-        char* variables[MAX_VARIABLES];                // Device array of pointer to device memory addresses for variable storage
-        unsigned int type_size[MAX_VARIABLES];         // Device array of the types of registered variables
-        unsigned int elements[MAX_VARIABLES];
-        unsigned int count[MAX_VARIABLES];
-    };
     /**
      * Main cuRVE variable hashing function
      *
@@ -61,6 +54,16 @@ class Curve {
     __host__ static VariableHash variableRuntimeHash(unsigned int num);
 #endif
 };
+/**
+ * Representation of Curve's hashtable
+ */
+struct CurveTable {
+    Curve::VariableHash hashes[Curve::MAX_VARIABLES];            // Device array of the hash values of registered variables
+    char* variables[Curve::MAX_VARIABLES];                // Device array of pointer to device memory addresses for variable storage
+    unsigned int type_size[Curve::MAX_VARIABLES];         // Device array of the types of registered variables
+    unsigned int elements[Curve::MAX_VARIABLES];
+    unsigned int count[Curve::MAX_VARIABLES];
+};
 
 /* TEMPLATE HASHING FUNCTIONS */