Place curve into flamegpu::detail::curve

CPPLINT.cfg

@@ -1,3 +1,3 @@
 set noparent 
-linelength=240
+linelength=320
 filter=-legal/copyright,-readability/todo,-runtime/references,-build/c++11

include/flamegpu/gpu/CUDAAgent.h

@@ -14,7 +14,7 @@
 #include "flamegpu/gpu/CUDAAgentStateList.h"
 #include "flamegpu/model/AgentFunctionData.cuh"
 #include "flamegpu/model/SubAgentData.h"
-#include "flamegpu/runtime/cuRVE/curve_rtc.h"
+#include "flamegpu/runtime/detail/curve/curve_rtc.h"
 #include "flamegpu/sim/AgentInterface.h"
 
 namespace flamegpu {
@@ -38,7 +38,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<CurveRTCHost>> CUDARTCHeaderMap;
+     typedef std::map<const std::string, std::unique_ptr<detail::curve::CurveRTCHost>> CUDARTCHeaderMap;
     /**
      * Element type of CUDARTCFuncMap
      */
@@ -218,7 +218,7 @@ class CUDAAgent : public AgentInterface {
      * @param function_name the name of the RTC agent function or the agent function name suffixed with condition (if it is a function condition)
      */
     const jitify::experimental::KernelInstantiation& getRTCInstantiation(const std::string &function_name) const;
-    CurveRTCHost &getRTCHeader(const std::string &function_name) const;
+    detail::curve::CurveRTCHost &getRTCHeader(const std::string &function_name) const;
     /**
      * Returns the CUDARTCFuncMap
      */

include/flamegpu/gpu/CUDAMessage.h

@@ -16,7 +16,11 @@ namespace flamegpu {
 class CUDAScatter;
 struct AgentFunctionData;
 struct MessageData;
+namespace detail {
+namespace curve {
 class Curve;
+}  // namespace curve
+}  // namespace detail
 class MsgSpecialisationHandler;
 class CUDAAgent;
 /**

include/flamegpu/gpu/CUDASimulation.h

@@ -9,7 +9,7 @@
 
 #include "flamegpu/exception/FLAMEGPUDeviceException.h"
 #include "flamegpu/sim/Simulation.h"
-
+#include "flamegpu/runtime/detail/curve/curve.h"
 #include "flamegpu/gpu/CUDAScatter.h"
 #include "flamegpu/gpu/CUDAEnsemble.h"
 #include "flamegpu/runtime/utility/RandomManager.cuh"
@@ -425,7 +425,7 @@ class CUDASimulation : public Simulation {
        * Curve instance used for variable mapping
        * @todo Is this necessary? CUDAAgent/CUDAMessage have their own copy
        */
-      Curve &curve;
+      detail::curve::Curve &curve;
       /**
        * Resizes device random array during step()
        */
@@ -444,7 +444,7 @@ class CUDASimulation : public Simulation {
        */
       exception::DeviceExceptionManager exception;
 #endif
-      Singletons(Curve &curve, EnvironmentManager &environment) : curve(curve), environment(environment) { }
+      Singletons(detail::curve::Curve &curve, EnvironmentManager &environment) : curve(curve), environment(environment) { }
     } * singletons;
     /**
      * Common method for adding this Model's data to env manager

include/flamegpu/model/AgentFunctionData.cuh

@@ -29,7 +29,7 @@ struct AgentFunctionData {
 
     /**
      * The cuda kernel entry point for executing the agent function
-     * @see void agent_function_wrapper(Curve::NamespaceHash, Curve::NamespaceHash, Curve::NamespaceHash, Curve::NamespaceHash, const int, const void *, const unsigned int, const unsigned int)
+     * @see void agent_function_wrapper(detail::curve::Curve::NamespaceHash, detail::curve::Curve::NamespaceHash, detail::curve::Curve::NamespaceHash, detail::curve::Curve::NamespaceHash, const int, const void *, const unsigned int, const unsigned int)
      */
     AgentFunctionWrapper *func;
     /**
@@ -75,7 +75,7 @@ struct AgentFunctionData {
     bool has_agent_death = false;
     /**
      * The cuda kernel entry point for executing the agent function condition
-     * @see void agent_function_condition_wrapper(Curve::NamespaceHash, Curve::NamespaceHash, const int, const unsigned int, const unsigned int)
+     * @see void agent_function_condition_wrapper(detail::curve::Curve::NamespaceHash, detail::curve::Curve::NamespaceHash, const int, const unsigned int, const unsigned int)
      */
     AgentFunctionConditionWrapper *condition;
     /**

include/flamegpu/model/AgentFunctionDescription.h

@@ -12,7 +12,7 @@
 #include "flamegpu/model/DependencyNode.h"
 #include "flamegpu/model/LayerDescription.h"
 #include "flamegpu/runtime/messaging/BruteForce.h"
-#include "flamegpu/runtime/cuRVE/curve_rtc.h"
+#include "flamegpu/runtime/detail/curve/curve_rtc.h"
 
 #ifdef _MSC_VER
 #pragma warning(push, 2)
@@ -312,8 +312,8 @@ template<typename AgentFunction>
 AgentFunctionDescription &AgentDescription::newFunction(const std::string &function_name, AgentFunction) {
     if (agent->functions.find(function_name) == agent->functions.end()) {
         AgentFunctionWrapper *f = AgentFunction::fnPtr();
-        std::string in_t = CurveRTCHost::demangle(AgentFunction::inType().name());
-        std::string out_t = CurveRTCHost::demangle(AgentFunction::outType().name());
+        std::string in_t = detail::curve::CurveRTCHost::demangle(AgentFunction::inType().name());
+        std::string out_t = detail::curve::CurveRTCHost::demangle(AgentFunction::outType().name());
         auto rtn = std::shared_ptr<AgentFunctionData>(new AgentFunctionData(this->agent->shared_from_this(), function_name, f, in_t, out_t));
         agent->functions.emplace(function_name, rtn);
         return *rtn->description;

include/flamegpu/runtime/AgentFunction.cuh

@@ -17,11 +17,11 @@ typedef void(AgentFunctionWrapper)(
 #if !defined(SEATBELTS) || SEATBELTS
     exception::DeviceExceptionBuffer *error_buffer,
 #endif
-    Curve::NamespaceHash instance_id_hash,
-    Curve::NamespaceHash agent_func_name_hash,
-    Curve::NamespaceHash messagename_inp_hash,
-    Curve::NamespaceHash messagename_outp_hash,
-    Curve::NamespaceHash agent_output_hash,
+    detail::curve::Curve::NamespaceHash instance_id_hash,
+    detail::curve::Curve::NamespaceHash agent_func_name_hash,
+    detail::curve::Curve::NamespaceHash messagename_inp_hash,
+    detail::curve::Curve::NamespaceHash messagename_outp_hash,
+    detail::curve::Curve::NamespaceHash agent_output_hash,
     id_t *d_agent_output_nextID,
     const unsigned int popNo,
     const void *in_messagelist_metadata,
@@ -57,11 +57,11 @@ __global__ void agent_function_wrapper(
 #if !defined(SEATBELTS) || SEATBELTS
     exception::DeviceExceptionBuffer *error_buffer,
 #endif
-    Curve::NamespaceHash instance_id_hash,
-    Curve::NamespaceHash agent_func_name_hash,
-    Curve::NamespaceHash messagename_inp_hash,
-    Curve::NamespaceHash messagename_outp_hash,
-    Curve::NamespaceHash agent_output_hash,
+    detail::curve::Curve::NamespaceHash instance_id_hash,
+    detail::curve::Curve::NamespaceHash agent_func_name_hash,
+    detail::curve::Curve::NamespaceHash messagename_inp_hash,
+    detail::curve::Curve::NamespaceHash messagename_outp_hash,
+    detail::curve::Curve::NamespaceHash agent_output_hash,
     id_t *d_agent_output_nextID,
     const unsigned int popNo,
     const void *in_messagelist_metadata,

include/flamegpu/runtime/AgentFunctionCondition.h

@@ -13,8 +13,8 @@ typedef void(AgentFunctionConditionWrapper)(
 #if !defined(SEATBELTS) || SEATBELTS
     exception::DeviceExceptionBuffer *error_buffer,
 #endif
-    Curve::NamespaceHash instance_id_hash,
-    Curve::NamespaceHash agent_func_name_hash,
+    detail::curve::Curve::NamespaceHash instance_id_hash,
+    detail::curve::Curve::NamespaceHash agent_func_name_hash,
     const unsigned int popNo,
     curandState *d_rng,
     unsigned int *scanFlag_conditionResult);  // Can't put __global__ in a typedef
@@ -36,8 +36,8 @@ __global__ void agent_function_condition_wrapper(
 #if !defined(SEATBELTS) || SEATBELTS
     exception::DeviceExceptionBuffer *error_buffer,
 #endif
-    Curve::NamespaceHash instance_id_hash,
-    Curve::NamespaceHash agent_func_name_hash,
+    detail::curve::Curve::NamespaceHash instance_id_hash,
+    detail::curve::Curve::NamespaceHash agent_func_name_hash,
     const unsigned int popNo,
     curandState *d_rng,
     unsigned int *scanFlag_conditionResult) {

include/flamegpu/runtime/DeviceAPI.h

@@ -7,7 +7,7 @@
 #include <limits>
 
 #ifndef __CUDACC_RTC__
-#include "flamegpu/runtime/cuRVE/curve.h"
+#include "flamegpu/runtime/detail/curve/curve.h"
 #else
 #include "dynamic/curve_rtc_dynamic.h"
 #endif  // !_RTC
@@ -32,8 +32,8 @@ class ReadOnlyDeviceAPI {
 #if !defined(SEATBELTS) || SEATBELTS
         exception::DeviceExceptionBuffer *error_buffer,
 #endif
-        Curve::NamespaceHash,
-        Curve::NamespaceHash,
+        detail::curve::Curve::NamespaceHash,
+        detail::curve::Curve::NamespaceHash,
         const unsigned int,
         curandState *,
         unsigned int *);
@@ -45,8 +45,8 @@ class ReadOnlyDeviceAPI {
      * @param d_rng Pointer to the device random state buffer to be used
      */
     __device__ ReadOnlyDeviceAPI(
-        const Curve::NamespaceHash &instance_id_hash,
-        const Curve::NamespaceHash &agentfuncname_hash,
+        const detail::curve::Curve::NamespaceHash &instance_id_hash,
+        const detail::curve::Curve::NamespaceHash &agentfuncname_hash,
         curandState *&d_rng)
         : random(AgentRandom(&d_rng[getThreadIndex()]))
         , environment(DeviceEnvironment(instance_id_hash))
@@ -123,7 +123,7 @@ class ReadOnlyDeviceAPI {
     }
 
  protected:
-    Curve::NamespaceHash agent_func_name_hash;
+    detail::curve::Curve::NamespaceHash agent_func_name_hash;
 };
 
 /** @brief    A flame gpu api class for the device runtime only
@@ -143,11 +143,11 @@ class DeviceAPI : public ReadOnlyDeviceAPI{
 #if !defined(SEATBELTS) || SEATBELTS
         exception::DeviceExceptionBuffer *error_buffer,
 #endif
-        Curve::NamespaceHash,
-        Curve::NamespaceHash,
-        Curve::NamespaceHash,
-        Curve::NamespaceHash,
-        Curve::NamespaceHash,
+        detail::curve::Curve::NamespaceHash,
+        detail::curve::Curve::NamespaceHash,
+        detail::curve::Curve::NamespaceHash,
+        detail::curve::Curve::NamespaceHash,
+        detail::curve::Curve::NamespaceHash,
         id_t*,
         const unsigned int,
         const void *,
@@ -169,7 +169,7 @@ class DeviceAPI : public ReadOnlyDeviceAPI{
          * @param d_agent_output_nextID Pointer to global memory holding the IDs to be assigned to new agents (selected via atomic inc)
          * @param scan_flag_agentOutput Pointer to (the start of) buffer of scan flags to be set true if this thread outputs an agent
          */
-        __device__ AgentOut(const Curve::NamespaceHash &aoh, id_t *&d_agent_output_nextID, unsigned int *&scan_flag_agentOutput)
+        __device__ AgentOut(const detail::curve::Curve::NamespaceHash &aoh, id_t *&d_agent_output_nextID, unsigned int *&scan_flag_agentOutput)
             : agent_output_hash(aoh)
             , scan_flag(scan_flag_agentOutput)
             , nextID(d_agent_output_nextID) { }
@@ -216,7 +216,7 @@ class DeviceAPI : public ReadOnlyDeviceAPI{
         /**
          * Curve hash used for accessing new agent variables
          */
-        const Curve::NamespaceHash agent_output_hash;
+        const detail::curve::Curve::NamespaceHash agent_output_hash;
         /**
          * Scan flag, defaults to 0, set to 1, to mark than agent is output
          */
@@ -243,9 +243,9 @@ class DeviceAPI : public ReadOnlyDeviceAPI{
      * @param msg_out Output message handler
      */
     __device__ DeviceAPI(
-        const Curve::NamespaceHash &instance_id_hash,
-        const Curve::NamespaceHash &agentfuncname_hash,
-        const Curve::NamespaceHash &_agent_output_hash,
+        const detail::curve::Curve::NamespaceHash &instance_id_hash,
+        const detail::curve::Curve::NamespaceHash &agentfuncname_hash,
+        const detail::curve::Curve::NamespaceHash &_agent_output_hash,
         id_t *&d_agent_output_nextID,
         curandState *&d_rng,
         unsigned int *&scanFlag_agentOutput,
@@ -305,7 +305,7 @@ __device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[N]) {
     const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x;
 
     // get the value from curve
-    T value = Curve::getAgentVariable<T>(variable_name, agent_func_name_hash , index);
+    T value = detail::curve::Curve::getAgentVariable<T>(variable_name, agent_func_name_hash , index);
 
     // return the variable from curve
     return value;
@@ -320,15 +320,15 @@ __device__ void DeviceAPI<MsgIn, MsgOut>::setVariable(const char(&variable_name)
     // simple indexing assumes index is the thread number (this may change later)
     const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x;
     // set the variable using curve
-    Curve::setAgentVariable<T>(variable_name, agent_func_name_hash,  value, index);
+    detail::curve::Curve::setAgentVariable<T>(variable_name, agent_func_name_hash,  value, index);
 }
 template<typename T, unsigned int N, unsigned int M>
 __device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[M], const unsigned int &array_index) {
     // simple indexing assumes index is the thread number (this may change later)
     const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x;
 
     // get the value from curve
-    T value = Curve::getAgentArrayVariable<T, N>(variable_name, agent_func_name_hash , index, array_index);
+    T value = detail::curve::Curve::getAgentArrayVariable<T, N>(variable_name, agent_func_name_hash , index, array_index);
 
     // return the variable from curve
     return value;
@@ -344,7 +344,7 @@ __device__ void DeviceAPI<MsgIn, MsgOut>::setVariable(const char(&variable_name)
     const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x;
 
     // set the variable using curve
-    Curve::setAgentArrayVariable<T, N>(variable_name , agent_func_name_hash,  value, index, array_index);
+    detail::curve::Curve::setAgentArrayVariable<T, N>(variable_name , agent_func_name_hash,  value, index, array_index);
 }
 
 template<typename MsgIn, typename MsgOut>
@@ -359,7 +359,7 @@ __device__ void DeviceAPI<MsgIn, MsgOut>::AgentOut::setVariable(const char(&vari
             const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x;
 
             // set the variable using curve
-            Curve::setNewAgentVariable<T>(variable_name, agent_output_hash, value, index);
+            detail::curve::Curve::setNewAgentVariable<T>(variable_name, agent_output_hash, value, index);
 
             // Mark scan flag
             genID();
@@ -381,7 +381,7 @@ __device__ void DeviceAPI<MsgIn, MsgOut>::AgentOut::setVariable(const char(&vari
         const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x;
 
         // set the variable using curve
-        Curve::setNewAgentArrayVariable<T, N>(variable_name, agent_output_hash, value, index, array_index);
+        detail::curve::Curve::setNewAgentArrayVariable<T, N>(variable_name, agent_output_hash, value, index, array_index);
 
         // Mark scan flag
         genID();
@@ -410,7 +410,7 @@ __device__ void DeviceAPI<MsgIn, MsgOut>::AgentOut::genID() const {
     if (this->id == ID_NOT_SET) {
         this->id = atomicInc(this->nextID, std::numeric_limits<id_t>().max());
         const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x;
-        Curve::setNewAgentVariable<id_t>("_id", agent_output_hash, this->id, index);  // Can't use ID_VARIABLE_NAME inline, as it isn't of char[N] type
+        detail::curve::Curve::setNewAgentVariable<id_t>("_id", agent_output_hash, this->id, index);  // Can't use ID_VARIABLE_NAME inline, as it isn't of char[N] type
         this->scan_flag[index] = 1;
     }
 }

include/flamegpu/runtime/cuRVE/curve.h → include/flamegpu/runtime/detail/curve/curve.h

@@ -1,5 +1,5 @@
-#ifndef INCLUDE_FLAMEGPU_RUNTIME_CURVE_CURVE_H_
-#define INCLUDE_FLAMEGPU_RUNTIME_CURVE_CURVE_H_
+#ifndef INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_H_
+#define INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_H_
 
 /**
  * @file   curve.h
@@ -24,6 +24,8 @@
 #include "flamegpu/exception/FLAMEGPUDeviceException.h"
 
 namespace flamegpu {
+namespace detail {
+namespace curve {
 
 /** @brief    A cuRVE instance.
  *
@@ -564,11 +566,10 @@ class Curve {
     /**
      * Has access to call purge
      */
-    friend class CUDASimulation;
+    friend class flamegpu::CUDASimulation;
     /**
      * Wipes out host mirrors of device memory
      * Only really to be used after calls to cudaDeviceReset()
-     * @note Only currently used after some tests
      */
     __host__ void purge();
 
@@ -595,14 +596,12 @@ class Curve {
 };
 
 
-namespace curve {
 namespace detail {
     extern __constant__ Curve::VariableHash d_hashes[Curve::MAX_VARIABLES];   // Device array of the hash values of registered variables
     extern __device__ char* d_variables[Curve::MAX_VARIABLES];                // Device array of pointer to device memory addresses for variable storage
     extern __constant__ size_t d_sizes[Curve::MAX_VARIABLES];                // Device array of the types of registered variables
     extern __constant__ unsigned int d_lengths[Curve::MAX_VARIABLES];
 }  // namespace detail
-}  // namespace curve
 
 
 /* TEMPLATE HASHING FUNCTIONS */
@@ -963,6 +962,8 @@ __device__ __forceinline__ void Curve::setArrayVariable(const char(&variableName
     return setArrayVariableByHash<T, N>(variable_hash + namespace_hash, variable, agent_index, array_index);
 }
 
+}  // namespace curve
+}  // namespace detail
 }  // namespace flamegpu
 
-#endif  // INCLUDE_FLAMEGPU_RUNTIME_CURVE_CURVE_H_
+#endif  // INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_H_

include/flamegpu/runtime/cuRVE/curve_rtc.h → include/flamegpu/runtime/detail/curve/curve_rtc.h

@@ -1,5 +1,5 @@
-#ifndef INCLUDE_FLAMEGPU_RUNTIME_CURVE_CURVE_RTC_H_
-#define INCLUDE_FLAMEGPU_RUNTIME_CURVE_CURVE_RTC_H_
+#ifndef INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_RTC_H_
+#define INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_RTC_H_
 
 #include <cstring>
 #include <string>
@@ -13,6 +13,8 @@ class KernelInstantiation;
 }  // namespace experimental
 }  // namespace jitify
 namespace flamegpu {
+namespace detail {
+namespace curve {
 
 
 /**
@@ -314,6 +316,9 @@ class CurveRTCHost {
      */
     std::map<std::string, RTCEnvVariableProperties> RTCEnvVariables;
 };
+
+}  // namespace curve
+}  // namespace detail
 }  // namespace flamegpu
 
-#endif  // INCLUDE_FLAMEGPU_RUNTIME_CURVE_CURVE_RTC_H_
+#endif  // INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_RTC_H_