diff --git a/CPPLINT.cfg b/CPPLINT.cfg index 50682c91b..ea2946437 100644 --- a/CPPLINT.cfg +++ b/CPPLINT.cfg @@ -1,3 +1,3 @@ set noparent -linelength=240 +linelength=320 filter=-legal/copyright,-readability/todo,-runtime/references,-build/c++11 diff --git a/include/flamegpu/gpu/CUDAAgent.h b/include/flamegpu/gpu/CUDAAgent.h index 79c2aa6a7..3a236d552 100644 --- a/include/flamegpu/gpu/CUDAAgent.h +++ b/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> CUDARTCFuncMap; - typedef std::map> CUDARTCHeaderMap; + typedef std::map> 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 */ diff --git a/include/flamegpu/gpu/CUDAMessage.h b/include/flamegpu/gpu/CUDAMessage.h index a62efceff..35b415f01 100644 --- a/include/flamegpu/gpu/CUDAMessage.h +++ b/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; /** diff --git a/include/flamegpu/gpu/CUDASimulation.h b/include/flamegpu/gpu/CUDASimulation.h index 3644e1741..f3852fba0 100644 --- a/include/flamegpu/gpu/CUDASimulation.h +++ b/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 diff --git a/include/flamegpu/model/AgentFunctionData.cuh b/include/flamegpu/model/AgentFunctionData.cuh index 6dba6991a..f4fa515d7 100644 --- a/include/flamegpu/model/AgentFunctionData.cuh +++ b/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; /** diff --git a/include/flamegpu/model/AgentFunctionDescription.h b/include/flamegpu/model/AgentFunctionDescription.h index c1878837b..7ac284dfb 100644 --- a/include/flamegpu/model/AgentFunctionDescription.h +++ b/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 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(new AgentFunctionData(this->agent->shared_from_this(), function_name, f, in_t, out_t)); agent->functions.emplace(function_name, rtn); return *rtn->description; diff --git a/include/flamegpu/runtime/AgentFunction.cuh b/include/flamegpu/runtime/AgentFunction.cuh index b8db2d8f4..d68e52fc9 100644 --- a/include/flamegpu/runtime/AgentFunction.cuh +++ b/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, diff --git a/include/flamegpu/runtime/AgentFunctionCondition.h b/include/flamegpu/runtime/AgentFunctionCondition.h index 38b920dbe..7ba7a0548 100644 --- a/include/flamegpu/runtime/AgentFunctionCondition.h +++ b/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) { diff --git a/include/flamegpu/runtime/DeviceAPI.h b/include/flamegpu/runtime/DeviceAPI.h index 72709d5d1..4a92ebb3a 100644 --- a/include/flamegpu/runtime/DeviceAPI.h +++ b/include/flamegpu/runtime/DeviceAPI.h @@ -7,7 +7,7 @@ #include #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(variable_name, agent_func_name_hash , index); + T value = detail::curve::Curve::getAgentVariable(variable_name, agent_func_name_hash , index); // return the variable from curve return value; @@ -320,7 +320,7 @@ __device__ void DeviceAPI::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(variable_name, agent_func_name_hash, value, index); + detail::curve::Curve::setAgentVariable(variable_name, agent_func_name_hash, value, index); } template __device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[M], const unsigned int &array_index) { @@ -328,7 +328,7 @@ __device__ T ReadOnlyDeviceAPI::getVariable(const char(&variable_name)[M], const const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // get the value from curve - T value = Curve::getAgentArrayVariable(variable_name, agent_func_name_hash , index, array_index); + T value = detail::curve::Curve::getAgentArrayVariable(variable_name, agent_func_name_hash , index, array_index); // return the variable from curve return value; @@ -344,7 +344,7 @@ __device__ void DeviceAPI::setVariable(const char(&variable_name) const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // set the variable using curve - Curve::setAgentArrayVariable(variable_name , agent_func_name_hash, value, index, array_index); + detail::curve::Curve::setAgentArrayVariable(variable_name , agent_func_name_hash, value, index, array_index); } template @@ -359,7 +359,7 @@ __device__ void DeviceAPI::AgentOut::setVariable(const char(&vari const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // set the variable using curve - Curve::setNewAgentVariable(variable_name, agent_output_hash, value, index); + detail::curve::Curve::setNewAgentVariable(variable_name, agent_output_hash, value, index); // Mark scan flag genID(); @@ -381,7 +381,7 @@ __device__ void DeviceAPI::AgentOut::setVariable(const char(&vari const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // set the variable using curve - Curve::setNewAgentArrayVariable(variable_name, agent_output_hash, value, index, array_index); + detail::curve::Curve::setNewAgentArrayVariable(variable_name, agent_output_hash, value, index, array_index); // Mark scan flag genID(); @@ -410,7 +410,7 @@ __device__ void DeviceAPI::AgentOut::genID() const { if (this->id == ID_NOT_SET) { this->id = atomicInc(this->nextID, std::numeric_limits().max()); const unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; - Curve::setNewAgentVariable("_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", 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; } } diff --git a/include/flamegpu/runtime/cuRVE/curve.h b/include/flamegpu/runtime/detail/curve/curve.h similarity index 99% rename from include/flamegpu/runtime/cuRVE/curve.h rename to include/flamegpu/runtime/detail/curve/curve.h index 8f8263b6f..47117bd5a 100644 --- a/include/flamegpu/runtime/cuRVE/curve.h +++ b/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(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_ diff --git a/include/flamegpu/runtime/cuRVE/curve_rtc.h b/include/flamegpu/runtime/detail/curve/curve_rtc.h similarity index 98% rename from include/flamegpu/runtime/cuRVE/curve_rtc.h rename to include/flamegpu/runtime/detail/curve/curve_rtc.h index 3b6f09143..c17e6b09d 100644 --- a/include/flamegpu/runtime/cuRVE/curve_rtc.h +++ b/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 #include @@ -13,6 +13,8 @@ class KernelInstantiation; } // namespace experimental } // namespace jitify namespace flamegpu { +namespace detail { +namespace curve { /** @@ -314,6 +316,9 @@ class CurveRTCHost { */ std::map RTCEnvVariables; }; + +} // namespace curve +} // namespace detail } // namespace flamegpu -#endif // INCLUDE_FLAMEGPU_RUNTIME_CURVE_CURVE_RTC_H_ +#endif // INCLUDE_FLAMEGPU_RUNTIME_DETAIL_CURVE_CURVE_RTC_H_ diff --git a/include/flamegpu/runtime/messaging/Array/ArrayDevice.h b/include/flamegpu/runtime/messaging/Array/ArrayDevice.h index 086bdcce0..6d3ec4e4c 100644 --- a/include/flamegpu/runtime/messaging/Array/ArrayDevice.h +++ b/include/flamegpu/runtime/messaging/Array/ArrayDevice.h @@ -205,7 +205,7 @@ class MsgArray::In { * @param x Search origin x coord * @param _radius Search radius */ - __device__ inline Filter(const size_type &_length, const Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &_radius); + __device__ inline Filter(const size_type &_length, const detail::curve::Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &_radius); /** * Returns an iterator to the start of the message list subset about the search origin */ @@ -239,7 +239,7 @@ class MsgArray::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; }; /** * Constructer @@ -248,7 +248,7 @@ class MsgArray::In { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param metadata Reinterpreted as type MsgArray::MetaData to extract length */ - __device__ In(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *metadata) + __device__ In(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *metadata) : combined_hash(agentfn_hash + msg_hash) , length(reinterpret_cast(metadata)->length) { } @@ -292,7 +292,7 @@ class MsgArray::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Metadata struct for accessing messages */ @@ -312,7 +312,7 @@ class MsgArray::Out { * @param _metadata Message specialisation specific metadata struct (of type MsgArray::MetaData) * @param scan_flag_messageOutput Scan flag array for optional message output */ - __device__ Out(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata, unsigned int *scan_flag_messageOutput) + __device__ Out(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata, unsigned int *scan_flag_messageOutput) : combined_hash(agentfn_hash + msg_hash) , scan_flag(scan_flag_messageOutput) #if !defined(SEATBELTS) || SEATBELTS @@ -341,7 +341,7 @@ class MsgArray::Out { * CURVE hash for accessing message data * agentfn_hash + msg_hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Scan flag array for optional message output */ @@ -362,7 +362,7 @@ __device__ T MsgArray::In::Message::getVariable(const char(&variable_name)[N]) c } #endif // get the value from curve using the stored hashes and message index. - return Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index); + return detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index); } template __device__ T MsgArray::In::Filter::Message::getVariable(const char(&variable_name)[N]) const { @@ -374,7 +374,7 @@ __device__ T MsgArray::In::Filter::Message::getVariable(const char(&variable_nam } #endif // get the value from curve using the stored hashes and message index. - return Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index_1d); + return detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index_1d); } template @@ -385,7 +385,7 @@ __device__ void MsgArray::Out::setVariable(const char(&variable_name)[N], T valu unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // set the variable using curve - Curve::setMessageVariable(variable_name, combined_hash, value, index); + detail::curve::Curve::setMessageVariable(variable_name, combined_hash, value, index); // setIndex() sets the optional msg scan flag } @@ -403,12 +403,12 @@ __device__ void MsgArray::Out::setIndex(const size_type &id) const { #endif // set the variable using curve - Curve::setMessageVariable("___INDEX", combined_hash, id, index); + detail::curve::Curve::setMessageVariable("___INDEX", combined_hash, id, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; } -__device__ MsgArray::In::Filter::Filter(const size_type &_length, const Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &_radius) +__device__ MsgArray::In::Filter::Filter(const size_type &_length, const detail::curve::Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &_radius) : radius(_radius) , length(_length) , combined_hash(_combined_hash) { diff --git a/include/flamegpu/runtime/messaging/Array2D/Array2DDevice.h b/include/flamegpu/runtime/messaging/Array2D/Array2DDevice.h index cd944d788..53dc12f47 100644 --- a/include/flamegpu/runtime/messaging/Array2D/Array2DDevice.h +++ b/include/flamegpu/runtime/messaging/Array2D/Array2DDevice.h @@ -211,7 +211,7 @@ class MsgArray2D::In { * @param y Search origin y coord * @param _radius Search radius */ - inline __device__ Filter(const MetaData *_metadata, const Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &y, const size_type &_radius); + inline __device__ Filter(const MetaData *_metadata, const detail::curve::Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &y, const size_type &_radius); /** * Returns an iterator to the start of the message list subset about the search origin */ @@ -245,7 +245,7 @@ class MsgArray2D::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; }; /** * Constructer @@ -254,7 +254,7 @@ class MsgArray2D::In { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param _metadata Reinterpreted as type MsgArray2D::MetaData */ - __device__ In(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata) + __device__ In(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata) : combined_hash(agentfn_hash + msg_hash) , metadata(reinterpret_cast(_metadata)) { } @@ -315,7 +315,7 @@ class MsgArray2D::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Metadata struct for accessing messages */ @@ -336,7 +336,7 @@ class MsgArray2D::Out { * @param _metadata Message specialisation specific metadata struct (of type MsgArray2D::MetaData) * @param scan_flag_messageOutput Scan flag array for optional message output */ - __device__ Out(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata, unsigned int *scan_flag_messageOutput) + __device__ Out(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata, unsigned int *scan_flag_messageOutput) : combined_hash(agentfn_hash + msg_hash) , scan_flag(scan_flag_messageOutput) , metadata(reinterpret_cast(_metadata)) @@ -361,7 +361,7 @@ class MsgArray2D::Out { * CURVE hash for accessing message data * agentfn_hash + msg_hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Scan flag array for optional message output */ @@ -384,7 +384,7 @@ __device__ T MsgArray2D::In::Message::getVariable(const char(&variable_name)[N]) } #endif // get the value from curve using the stored hashes and message index. - return Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index); + return detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index); } template __device__ T MsgArray2D::In::Filter::Message::getVariable(const char(&variable_name)[N]) const { @@ -396,7 +396,7 @@ __device__ T MsgArray2D::In::Filter::Message::getVariable(const char(&variable_n } #endif // get the value from curve using the stored hashes and message index. - return Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index_1d); + return detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index_1d); } template @@ -407,7 +407,7 @@ __device__ void MsgArray2D::Out::setVariable(const char(&variable_name)[N], T va unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // set the variable using curve - Curve::setMessageVariable(variable_name, combined_hash, value, index); + detail::curve::Curve::setMessageVariable(variable_name, combined_hash, value, index); // setIndex() sets the optional msg scan flag } @@ -428,12 +428,12 @@ __device__ void MsgArray2D::Out::setIndex(const size_type &x, const size_type &y } #endif // set the variable using curve - Curve::setMessageVariable("___INDEX", combined_hash, index_1d, index); + detail::curve::Curve::setMessageVariable("___INDEX", combined_hash, index_1d, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; } -__device__ MsgArray2D::In::Filter::Filter(const MetaData *_metadata, const Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &y, const size_type &_radius) +__device__ MsgArray2D::In::Filter::Filter(const MetaData *_metadata, const detail::curve::Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &y, const size_type &_radius) : radius(_radius) , metadata(_metadata) , combined_hash(_combined_hash) { diff --git a/include/flamegpu/runtime/messaging/Array3D/Array3DDevice.h b/include/flamegpu/runtime/messaging/Array3D/Array3DDevice.h index 297089699..3ad7d2efd 100644 --- a/include/flamegpu/runtime/messaging/Array3D/Array3DDevice.h +++ b/include/flamegpu/runtime/messaging/Array3D/Array3DDevice.h @@ -225,7 +225,7 @@ class MsgArray3D::In { * @param z Search origin z coord * @param _radius Search radius */ - inline __device__ Filter(const MetaData *_metadata, const Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &y, const size_type &z, const size_type &_radius); + inline __device__ Filter(const MetaData *_metadata, const detail::curve::Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &y, const size_type &z, const size_type &_radius); /** * Returns an iterator to the start of the message list subset about the search origin */ @@ -259,7 +259,7 @@ class MsgArray3D::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; }; /** * Constructer @@ -268,7 +268,7 @@ class MsgArray3D::In { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param _metadata Reinterpreted as type MsgArray3D::MetaData */ - __device__ In(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata) + __device__ In(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata) : combined_hash(agentfn_hash + msg_hash) , metadata(reinterpret_cast(_metadata)) { } @@ -337,7 +337,7 @@ class MsgArray3D::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Metadata struct for accessing messages */ @@ -358,7 +358,7 @@ class MsgArray3D::Out { * @param _metadata Message specialisation specific metadata struct (of type MsgArray3D::MetaData) * @param scan_flag_messageOutput Scan flag array for optional message output */ - __device__ Out(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata, unsigned int *scan_flag_messageOutput) + __device__ Out(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata, unsigned int *scan_flag_messageOutput) : combined_hash(agentfn_hash + msg_hash) , scan_flag(scan_flag_messageOutput) , metadata(reinterpret_cast(_metadata)) @@ -383,7 +383,7 @@ class MsgArray3D::Out { * CURVE hash for accessing message data * agentfn_hash + msg_hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Scan flag array for optional message output */ @@ -404,7 +404,7 @@ __device__ T MsgArray3D::In::Message::getVariable(const char(&variable_name)[N]) } #endif // get the value from curve using the stored hashes and message index. - return Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index); + return detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index); } template __device__ T MsgArray3D::In::Filter::Message::getVariable(const char(&variable_name)[N]) const { @@ -416,7 +416,7 @@ __device__ T MsgArray3D::In::Filter::Message::getVariable(const char(&variable_n } #endif // get the value from curve using the stored hashes and message index. - return Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index_1d); + return detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, index_1d); } template @@ -427,7 +427,7 @@ __device__ void MsgArray3D::Out::setVariable(const char(&variable_name)[N], T va unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // set the variable using curve - Curve::setMessageVariable(variable_name, combined_hash, value, index); + detail::curve::Curve::setMessageVariable(variable_name, combined_hash, value, index); // setIndex() sets the optional msg scan flag } @@ -450,12 +450,12 @@ __device__ inline void MsgArray3D::Out::setIndex(const size_type &x, const size_ #endif // set the variable using curve - Curve::setMessageVariable("___INDEX", combined_hash, index_1d, index); + detail::curve::Curve::setMessageVariable("___INDEX", combined_hash, index_1d, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; } -__device__ inline MsgArray3D::In::Filter::Filter(const MetaData *_metadata, const Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &y, const size_type &z, const size_type &_radius) +__device__ inline MsgArray3D::In::Filter::Filter(const MetaData *_metadata, const detail::curve::Curve::NamespaceHash &_combined_hash, const size_type &x, const size_type &y, const size_type &z, const size_type &_radius) : radius(_radius) , metadata(_metadata) , combined_hash(_combined_hash) { diff --git a/include/flamegpu/runtime/messaging/BruteForce/BruteForceDevice.h b/include/flamegpu/runtime/messaging/BruteForce/BruteForceDevice.h index a1fda3cff..de8130463 100644 --- a/include/flamegpu/runtime/messaging/BruteForce/BruteForceDevice.h +++ b/include/flamegpu/runtime/messaging/BruteForce/BruteForceDevice.h @@ -29,7 +29,7 @@ class MsgBruteForce::In { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param metadata Reinterpreted as type MsgBruteForce::MetaData to extract length */ - __device__ In(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *metadata) + __device__ In(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *metadata) : combined_hash(agentfn_hash + msg_hash) , len(reinterpret_cast(metadata)->length) { } @@ -153,7 +153,7 @@ class MsgBruteForce::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Total number of messages in the message list */ @@ -175,7 +175,7 @@ class MsgBruteForce::Out { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param scan_flag_messageOutput Scan flag array for optional message output */ - __device__ Out(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *, unsigned int *scan_flag_messageOutput) + __device__ Out(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *, unsigned int *scan_flag_messageOutput) : combined_hash(agentfn_hash + msg_hash) , scan_flag(scan_flag_messageOutput) { } @@ -195,7 +195,7 @@ class MsgBruteForce::Out { * CURVE hash for accessing message data * agentfn_hash + msg_hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Scan flag array for optional message output */ @@ -212,7 +212,7 @@ __device__ T MsgBruteForce::In::Message::getVariable(const char(&variable_name)[ } #endif // get the value from curve using the stored hashes and message index. - T value = Curve::getMessageVariable_ldg(variable_name, this->_parent.combined_hash, index); + T value = detail::curve::Curve::getMessageVariable_ldg(variable_name, this->_parent.combined_hash, index); return value; } @@ -226,7 +226,7 @@ __device__ void MsgBruteForce::Out::setVariable(const char(&variable_name)[N], T // Todo: checking if the output message type is single or optional? (d_message_type) // set the variable using curve - Curve::setMessageVariable(variable_name, combined_hash, value, index); + detail::curve::Curve::setMessageVariable(variable_name, combined_hash, value, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; diff --git a/include/flamegpu/runtime/messaging/BruteForce/BruteForceHost.h b/include/flamegpu/runtime/messaging/BruteForce/BruteForceHost.h index 9057821cf..d3c3a859d 100644 --- a/include/flamegpu/runtime/messaging/BruteForce/BruteForceHost.h +++ b/include/flamegpu/runtime/messaging/BruteForce/BruteForceHost.h @@ -11,7 +11,7 @@ #include "flamegpu/model/Variable.h" #include "flamegpu/gpu/detail/CUDAErrorChecking.cuh" -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #include "flamegpu/runtime/messaging/None/NoneHost.h" #include "flamegpu/runtime/messaging/BruteForce.h" diff --git a/include/flamegpu/runtime/messaging/Bucket.h b/include/flamegpu/runtime/messaging/Bucket.h index 051f5b4a6..f2561e1e3 100644 --- a/include/flamegpu/runtime/messaging/Bucket.h +++ b/include/flamegpu/runtime/messaging/Bucket.h @@ -5,7 +5,7 @@ #include #include -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #endif // __CUDACC_RTC__ #include "flamegpu/runtime/messaging/None.h" diff --git a/include/flamegpu/runtime/messaging/Bucket/BucketDevice.h b/include/flamegpu/runtime/messaging/Bucket/BucketDevice.h index 2493dc08b..3cdafe8d6 100644 --- a/include/flamegpu/runtime/messaging/Bucket/BucketDevice.h +++ b/include/flamegpu/runtime/messaging/Bucket/BucketDevice.h @@ -139,7 +139,7 @@ class MsgBucket::In { * @param beginKey Inclusive first bucket of range to access * @param endKey Exclusive final bucket of range to access, this is the final bucket + 1 */ - inline __device__ Filter(const MetaData *_metadata, const Curve::NamespaceHash &combined_hash, const IntT &beginKey, const IntT &endKey); + inline __device__ Filter(const MetaData *_metadata, const detail::curve::Curve::NamespaceHash &combined_hash, const IntT &beginKey, const IntT &endKey); /** * Returns an iterator to the start of the message list subset about the search origin */ @@ -175,7 +175,7 @@ class MsgBucket::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; }; /** * Constructor @@ -184,7 +184,7 @@ class MsgBucket::In { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param _metadata Reinterpreted as type MsgBucket::MetaData */ - __device__ In(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata) + __device__ In(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata) : combined_hash(agentfn_hash + msg_hash) , metadata(reinterpret_cast(_metadata)) { } @@ -233,7 +233,7 @@ class MsgBucket::In { * CURVE hash for accessing message data * agentfn_hash + msg_hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Device pointer to metadata required for accessing data structure * e.g. PBM, search origin, environment bounds @@ -255,7 +255,7 @@ class MsgBucket::Out : public MsgBruteForce::Out { * @param _metadata Message specialisation specific metadata struct (of type MsgBucket::MetaData) * @param scan_flag_messageOutput Scan flag array for optional message output */ - __device__ Out(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata, unsigned int *scan_flag_messageOutput) + __device__ Out(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata, unsigned int *scan_flag_messageOutput) : MsgBruteForce::Out(agentfn_hash, msg_hash, nullptr, scan_flag_messageOutput) #if !defined(SEATBELTS) || SEATBELTS , metadata(reinterpret_cast(_metadata)) @@ -275,7 +275,7 @@ class MsgBucket::Out : public MsgBruteForce::Out { const MetaData * const metadata; }; -__device__ MsgBucket::In::Filter::Filter(const MetaData* _metadata, const Curve::NamespaceHash &_combined_hash, const IntT& beginKey, const IntT& endKey) +__device__ MsgBucket::In::Filter::Filter(const MetaData* _metadata, const detail::curve::Curve::NamespaceHash &_combined_hash, const IntT& beginKey, const IntT& endKey) : bucket_begin(0) , bucket_end(0) , metadata(_metadata) @@ -296,7 +296,7 @@ __device__ void MsgBucket::Out::setKey(const IntT &key) const { } #endif // set the variables using curve - Curve::setMessageVariable("_key", combined_hash, key, index); + detail::curve::Curve::setMessageVariable("_key", combined_hash, key, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; @@ -312,7 +312,7 @@ __device__ T MsgBucket::In::Filter::Message::getVariable(const char(&variable_na } #endif // get the value from curve using the stored hashes and message index. - T value = Curve::getMessageVariable(variable_name, this->_parent.combined_hash, cell_index); + T value = detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, cell_index); return value; } } // namespace flamegpu diff --git a/include/flamegpu/runtime/messaging/None/NoneDevice.h b/include/flamegpu/runtime/messaging/None/NoneDevice.h index f9f747ba8..7e19fae78 100644 --- a/include/flamegpu/runtime/messaging/None/NoneDevice.h +++ b/include/flamegpu/runtime/messaging/None/NoneDevice.h @@ -2,7 +2,7 @@ #define INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_NONE_NONEDEVICE_H_ #ifndef __CUDACC_RTC__ -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #endif // __CUDACC_RTC__ #include "flamegpu/runtime/messaging/None.h" @@ -20,7 +20,7 @@ class MsgNone::In { * Requires CURVE hashes for agent function and message name to retrieve variable memory locations * Takes a device pointer to a struct for metadata related to accessing the messages (e.g. an index data structure) */ - __device__ In(Curve::NamespaceHash /*agent fn hash*/, Curve::NamespaceHash /*message name hash*/, const void * /*metadata*/) { + __device__ In(detail::curve::Curve::NamespaceHash /*agent fn hash*/, detail::curve::Curve::NamespaceHash /*message name hash*/, const void * /*metadata*/) { } }; /** @@ -34,7 +34,7 @@ class MsgNone::Out { * Requires CURVE hashes for agent function and message name to retrieve variable memory locations * Takes a device pointer to a struct for metadata related to accessing the messages (e.g. an index data structure) */ - __device__ Out(Curve::NamespaceHash /*agent fn hash*/, Curve::NamespaceHash /*message name hash*/, const void * /*metadata*/, unsigned int * /*scan_flag_messageOutput*/){ + __device__ Out(detail::curve::Curve::NamespaceHash /*agent fn hash*/, detail::curve::Curve::NamespaceHash /*message name hash*/, const void * /*metadata*/, unsigned int * /*scan_flag_messageOutput*/){ } }; diff --git a/include/flamegpu/runtime/messaging/None/NoneHost.h b/include/flamegpu/runtime/messaging/None/NoneHost.h index 1550ea9af..bc407e48c 100644 --- a/include/flamegpu/runtime/messaging/None/NoneHost.h +++ b/include/flamegpu/runtime/messaging/None/NoneHost.h @@ -1,7 +1,7 @@ #ifndef INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_NONE_NONEHOST_H_ #define INCLUDE_FLAMEGPU_RUNTIME_MESSAGING_NONE_NONEHOST_H_ -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #include "flamegpu/runtime/messaging/None.h" namespace flamegpu { diff --git a/include/flamegpu/runtime/messaging/Spatial2D/Spatial2DDevice.h b/include/flamegpu/runtime/messaging/Spatial2D/Spatial2DDevice.h index 68fdf9cc4..eaee38a0d 100644 --- a/include/flamegpu/runtime/messaging/Spatial2D/Spatial2DDevice.h +++ b/include/flamegpu/runtime/messaging/Spatial2D/Spatial2DDevice.h @@ -171,7 +171,7 @@ class MsgSpatial2D::In { * @param x Search origin x coord * @param y Search origin y coord */ - __device__ Filter(const MetaData *_metadata, const Curve::NamespaceHash &combined_hash, const float &x, const float &y); + __device__ Filter(const MetaData *_metadata, const detail::curve::Curve::NamespaceHash &combined_hash, const float &x, const float &y); /** * Returns an iterator to the start of the message list subset about the search origin */ @@ -206,7 +206,7 @@ class MsgSpatial2D::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; }; /** * Constructer @@ -215,7 +215,7 @@ class MsgSpatial2D::In { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param _metadata Reinterpreted as type MsgSpatial3D::MetaData */ - __device__ In(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata) + __device__ In(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata) : combined_hash(agentfn_hash + msg_hash) , metadata(reinterpret_cast(_metadata)) { } @@ -242,7 +242,7 @@ class MsgSpatial2D::In { * CURVE hash for accessing message data * agentfn_hash + msg_hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Device pointer to metadata required for accessing data structure * e.g. PBM, search origin, environment bounds @@ -263,7 +263,7 @@ class MsgSpatial2D::Out : public MsgBruteForce::Out { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param scan_flag_messageOutput Scan flag array for optional message output */ - __device__ Out(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *, unsigned int *scan_flag_messageOutput) + __device__ Out(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *, unsigned int *scan_flag_messageOutput) : MsgBruteForce::Out(agentfn_hash, msg_hash, nullptr, scan_flag_messageOutput) { } /** @@ -285,7 +285,7 @@ __device__ T MsgSpatial2D::In::Filter::Message::getVariable(const char(&variable } #endif // get the value from curve using the stored hashes and message index. - T value = Curve::getMessageVariable(variable_name, this->_parent.combined_hash, cell_index); + T value = detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, cell_index); return value; } @@ -318,14 +318,14 @@ __device__ inline void MsgSpatial2D::Out::setLocation(const float &x, const floa unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // + d_message_count; // set the variables using curve - Curve::setMessageVariable("x", combined_hash, x, index); - Curve::setMessageVariable("y", combined_hash, y, index); + detail::curve::Curve::setMessageVariable("x", combined_hash, x, index); + detail::curve::Curve::setMessageVariable("y", combined_hash, y, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; } -__device__ inline MsgSpatial2D::In::Filter::Filter(const MetaData* _metadata, const Curve::NamespaceHash &_combined_hash, const float& x, const float& y) +__device__ inline MsgSpatial2D::In::Filter::Filter(const MetaData* _metadata, const detail::curve::Curve::NamespaceHash &_combined_hash, const float& x, const float& y) : metadata(_metadata) , combined_hash(_combined_hash) { loc[0] = x; diff --git a/include/flamegpu/runtime/messaging/Spatial2D/Spatial2DHost.h b/include/flamegpu/runtime/messaging/Spatial2D/Spatial2DHost.h index ac5bda8f9..c319b2375 100644 --- a/include/flamegpu/runtime/messaging/Spatial2D/Spatial2DHost.h +++ b/include/flamegpu/runtime/messaging/Spatial2D/Spatial2DHost.h @@ -4,7 +4,7 @@ #include #include -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #include "flamegpu/model/Variable.h" #include "flamegpu/runtime/messaging/Spatial2D.h" #include "flamegpu/runtime/messaging/BruteForce/BruteForceHost.h" diff --git a/include/flamegpu/runtime/messaging/Spatial3D/Spatial3DDevice.h b/include/flamegpu/runtime/messaging/Spatial3D/Spatial3DDevice.h index 9d5709099..19105edd1 100644 --- a/include/flamegpu/runtime/messaging/Spatial3D/Spatial3DDevice.h +++ b/include/flamegpu/runtime/messaging/Spatial3D/Spatial3DDevice.h @@ -180,7 +180,7 @@ class MsgSpatial3D::In { * @param y Search origin y coord * @param z search origin z coord */ - __device__ Filter(const MetaData *_metadata, const Curve::NamespaceHash &combined_hash, const float &x, const float &y, const float &z); + __device__ Filter(const MetaData *_metadata, const detail::curve::Curve::NamespaceHash &combined_hash, const float &x, const float &y, const float &z); /** * Returns an iterator to the start of the message list subset about the search origin */ @@ -215,7 +215,7 @@ class MsgSpatial3D::In { * CURVE hash for accessing message data * agent function hash + message hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; }; /** @@ -225,7 +225,7 @@ class MsgSpatial3D::In { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param _metadata Reinterpreted as type MsgSpatial3D::MetaData */ - __device__ In(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *_metadata) + __device__ In(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *_metadata) : combined_hash(agentfn_hash + msg_hash) , metadata(reinterpret_cast(_metadata)) { } @@ -253,7 +253,7 @@ class MsgSpatial3D::In { * CURVE hash for accessing message data * agentfn_hash + msg_hash */ - Curve::NamespaceHash combined_hash; + detail::curve::Curve::NamespaceHash combined_hash; /** * Device pointer to metadata required for accessing data structure * e.g. PBM, search origin, environment bounds @@ -274,7 +274,7 @@ class MsgSpatial3D::Out : public MsgBruteForce::Out { * @param msg_hash Added to agentfn_hash to produce combined_hash * @param scan_flag_messageOutput Scan flag array for optional message output */ - __device__ Out(Curve::NamespaceHash agentfn_hash, Curve::NamespaceHash msg_hash, const void *, unsigned int *scan_flag_messageOutput) + __device__ Out(detail::curve::Curve::NamespaceHash agentfn_hash, detail::curve::Curve::NamespaceHash msg_hash, const void *, unsigned int *scan_flag_messageOutput) : MsgBruteForce::Out(agentfn_hash, msg_hash, nullptr, scan_flag_messageOutput) { } /** @@ -297,7 +297,7 @@ __device__ T MsgSpatial3D::In::Filter::Message::getVariable(const char(&variable } #endif // get the value from curve using the stored hashes and message index. - T value = Curve::getMessageVariable(variable_name, this->_parent.combined_hash, cell_index); + T value = detail::curve::Curve::getMessageVariable(variable_name, this->_parent.combined_hash, cell_index); return value; } @@ -334,15 +334,15 @@ __device__ inline void MsgSpatial3D::Out::setLocation(const float &x, const floa unsigned int index = (blockDim.x * blockIdx.x) + threadIdx.x; // + d_message_count; // set the variables using curve - Curve::setMessageVariable("x", combined_hash, x, index); - Curve::setMessageVariable("y", combined_hash, y, index); - Curve::setMessageVariable("z", combined_hash, z, index); + detail::curve::Curve::setMessageVariable("x", combined_hash, x, index); + detail::curve::Curve::setMessageVariable("y", combined_hash, y, index); + detail::curve::Curve::setMessageVariable("z", combined_hash, z, index); // Set scan flag incase the message is optional this->scan_flag[index] = 1; } -__device__ inline MsgSpatial3D::In::Filter::Filter(const MetaData* _metadata, const Curve::NamespaceHash &_combined_hash, const float& x, const float& y, const float& z) +__device__ inline MsgSpatial3D::In::Filter::Filter(const MetaData* _metadata, const detail::curve::Curve::NamespaceHash &_combined_hash, const float& x, const float& y, const float& z) : metadata(_metadata) , combined_hash(_combined_hash) { loc[0] = x; diff --git a/include/flamegpu/runtime/messaging/Spatial3D/Spatial3DHost.h b/include/flamegpu/runtime/messaging/Spatial3D/Spatial3DHost.h index 3d6e865d1..94836390c 100644 --- a/include/flamegpu/runtime/messaging/Spatial3D/Spatial3DHost.h +++ b/include/flamegpu/runtime/messaging/Spatial3D/Spatial3DHost.h @@ -5,7 +5,7 @@ #include #include "flamegpu/gpu/CUDAMessage.h" -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #include "flamegpu/util/nvtx.h" #include "flamegpu/runtime/messaging/Spatial3D.h" #include "flamegpu/runtime/messaging/Spatial2D/Spatial2DHost.h" diff --git a/include/flamegpu/runtime/utility/DeviceEnvironment.cuh b/include/flamegpu/runtime/utility/DeviceEnvironment.cuh index 20a79aabd..c1d213237 100644 --- a/include/flamegpu/runtime/utility/DeviceEnvironment.cuh +++ b/include/flamegpu/runtime/utility/DeviceEnvironment.cuh @@ -39,12 +39,12 @@ class DeviceEnvironment { /** * Hash of the model's name, this is added to CURVE_NAMESPACE_HASH and variable name hash to find curve hash */ - const Curve::NamespaceHash &modelname_hash; + const detail::curve::Curve::NamespaceHash &modelname_hash; /** * Constructor, requires the model name hash to init modelname_hash * @param _modelname_hash Hash of model name generated by curveGetVariableHash() */ - __device__ __forceinline__ DeviceEnvironment(const Curve::NamespaceHash &_modelname_hash) + __device__ __forceinline__ DeviceEnvironment(const detail::curve::Curve::NamespaceHash &_modelname_hash) : modelname_hash(_modelname_hash) { } public: @@ -87,38 +87,38 @@ class DeviceEnvironment { */ template __device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N]) const { - Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + Curve::variableHash(name); - const auto cv = Curve::getVariable(cvh); + detail::curve::Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + detail::curve::Curve::variableHash(name); + const auto cv = detail::curve::Curve::getVariable(cvh); #if !defined(SEATBELTS) || SEATBELTS - if (cv == Curve::UNKNOWN_VARIABLE) { + if (cv == detail::curve::Curve::UNKNOWN_VARIABLE) { DTHROW("Environment property with name: %s was not found.\n", name); - } else if (curve::detail::d_sizes[cv] != sizeof(T)) { - DTHROW("Environment property with name: %s type size mismatch %llu != %llu.\n", name, curve::detail::d_sizes[cv], sizeof(T)); + } else if (detail::curve::detail::d_sizes[cv] != sizeof(T)) { + DTHROW("Environment property with name: %s type size mismatch %llu != %llu.\n", name, detail::curve::detail::d_sizes[cv], sizeof(T)); } else { - return *reinterpret_cast(detail::c_envPropBuffer + reinterpret_cast(curve::detail::d_variables[cv])); + return *reinterpret_cast(detail::c_envPropBuffer + reinterpret_cast(detail::curve::detail::d_variables[cv])); } return {}; #else - return *reinterpret_cast(detail::c_envPropBuffer + reinterpret_cast(curve::detail::d_variables[cv])); + return *reinterpret_cast(detail::c_envPropBuffer + reinterpret_cast(detail::curve::detail::d_variables[cv])); #endif } template __device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N], const unsigned int &index) const { - Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + Curve::variableHash(name); - const auto cv = Curve::getVariable(cvh); + detail::curve::Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + detail::curve::Curve::variableHash(name); + const auto cv = detail::curve::Curve::getVariable(cvh); #if !defined(SEATBELTS) || SEATBELTS - if (cv == Curve::UNKNOWN_VARIABLE) { + if (cv == detail::curve::Curve::UNKNOWN_VARIABLE) { DTHROW("Environment property array with name: %s was not found.\n", name); - } else if (curve::detail::d_sizes[cv] != sizeof(T)) { - DTHROW("Environment property array with name: %s type size mismatch %llu != %llu.\n", name, curve::detail::d_sizes[cv], sizeof(T)); - } else if (curve::detail::d_lengths[cv] <= index) { - DTHROW("Environment property array with name: %s index %u is out of bounds (length %u).\n", name, index, curve::detail::d_lengths[cv]); + } else if (detail::curve::detail::d_sizes[cv] != sizeof(T)) { + DTHROW("Environment property array with name: %s type size mismatch %llu != %llu.\n", name, detail::curve::detail::d_sizes[cv], sizeof(T)); + } else if (detail::curve::detail::d_lengths[cv] <= index) { + DTHROW("Environment property array with name: %s index %u is out of bounds (length %u).\n", name, index, detail::curve::detail::d_lengths[cv]); } else { - return *(reinterpret_cast(detail::c_envPropBuffer + reinterpret_cast(curve::detail::d_variables[cv])) + index); + return *(reinterpret_cast(detail::c_envPropBuffer + reinterpret_cast(detail::curve::detail::d_variables[cv])) + index); } return {}; #else - return *(reinterpret_cast(detail::c_envPropBuffer + reinterpret_cast(curve::detail::d_variables[cv])) + index); + return *(reinterpret_cast(detail::c_envPropBuffer + reinterpret_cast(detail::curve::detail::d_variables[cv])) + index); #endif } @@ -127,8 +127,8 @@ __device__ __forceinline__ T DeviceEnvironment::getProperty(const char(&name)[N] */ template __device__ __forceinline__ bool DeviceEnvironment::containsProperty(const char(&name)[N]) const { - Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + Curve::variableHash(name); - return Curve::getVariable(cvh) != Curve::UNKNOWN_VARIABLE; + detail::curve::Curve::VariableHash cvh = CURVE_NAMESPACE_HASH() + modelname_hash + detail::curve::Curve::variableHash(name); + return detail::curve::Curve::getVariable(cvh) != detail::curve::Curve::UNKNOWN_VARIABLE; } #endif // __CUDACC_RTC__ diff --git a/include/flamegpu/runtime/utility/EnvironmentManager.cuh b/include/flamegpu/runtime/utility/EnvironmentManager.cuh index f4572828b..8afc4152f 100644 --- a/include/flamegpu/runtime/utility/EnvironmentManager.cuh +++ b/include/flamegpu/runtime/utility/EnvironmentManager.cuh @@ -21,7 +21,7 @@ #include "flamegpu/exception/FLAMEGPUException.h" #include "flamegpu/gpu/detail/CUDAErrorChecking.cuh" -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #include "flamegpu/util/Any.h" namespace flamegpu { @@ -251,7 +251,7 @@ class EnvironmentManager { * @param curve The Curve singleton instance to use, it is important that we purge curve for the correct device * @param instance_id instance_id of the CUDASimulation instance the properties are attached to */ - void free(Curve &curve, const unsigned int &instance_id); + void free(detail::curve::Curve &curve, const unsigned int &instance_id); /** * Adds a new environment property * @param name name used for accessing the property @@ -607,7 +607,7 @@ class EnvironmentManager { * Hash never changes, so we store a copy at creation * Also ensure the device constexpr version matches */ - const Curve::NamespaceHash CURVE_NAMESPACE_HASH; + const detail::curve::Curve::NamespaceHash CURVE_NAMESPACE_HASH; /** * Returns read-only access to the properties map * @note You must acquire a lock on mutex before calling this method @@ -648,7 +648,7 @@ class EnvironmentManager { * @param name Pair of the two items to produce the curve value hash * @note Not static, because eventually we might need to use curve singleton */ - Curve::VariableHash toHash(const NamePair &name) const; + detail::curve::Curve::VariableHash toHash(const NamePair &name) const; /** * Common add handler */ @@ -661,7 +661,7 @@ class EnvironmentManager { * @param isPureRTC If true, Curve collision warnings (debug build only) will be suppressed as they are irrelevant to RTC models * @note any EnvPROP */ - void defragment(Curve &curve, const DefragMap * mergeProps = nullptr, std::set newmaps = {}, bool isPureRTC = false); + void defragment(detail::curve::Curve &curve, const DefragMap * mergeProps = nullptr, std::set newmaps = {}, bool isPureRTC = false); /** * This is the RTC version of defragment() * RTC Constant offsets are fixed at RTC time, and exist in their own constant block. diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index f1c44ab80..d9aa05279 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -139,8 +139,8 @@ SET(SRC_INCLUDE ${FLAMEGPU_ROOT}/include/flamegpu/runtime/HostAPI_macros.h ${FLAMEGPU_ROOT}/include/flamegpu/runtime/HostAgentAPI.h ${FLAMEGPU_ROOT}/include/flamegpu/runtime/HostNewAgentAPI.h - ${FLAMEGPU_ROOT}/include/flamegpu/runtime/cuRVE/curve.h - ${FLAMEGPU_ROOT}/include/flamegpu/runtime/cuRVE/curve_rtc.h + ${FLAMEGPU_ROOT}/include/flamegpu/runtime/detail/curve/curve.h + ${FLAMEGPU_ROOT}/include/flamegpu/runtime/detail/curve/curve_rtc.h ${FLAMEGPU_ROOT}/include/flamegpu/runtime/messaging.h ${FLAMEGPU_ROOT}/include/flamegpu/runtime/messaging_device.h ${FLAMEGPU_ROOT}/include/flamegpu/runtime/messaging/None.h @@ -229,8 +229,8 @@ SET(SRC_FLAMEGPU2 ${FLAMEGPU_ROOT}/src/flamegpu/sim/SimRunner.cu ${FLAMEGPU_ROOT}/src/flamegpu/sim/SimLogger.cu ${FLAMEGPU_ROOT}/src/flamegpu/sim/Simulation.cu - ${FLAMEGPU_ROOT}/src/flamegpu/runtime/cuRVE/curve.cu - ${FLAMEGPU_ROOT}/src/flamegpu/runtime/cuRVE/curve_rtc.cpp + ${FLAMEGPU_ROOT}/src/flamegpu/runtime/detail/curve/curve.cu + ${FLAMEGPU_ROOT}/src/flamegpu/runtime/detail/curve/curve_rtc.cpp ${FLAMEGPU_ROOT}/src/flamegpu/runtime/HostAPI.cu ${FLAMEGPU_ROOT}/src/flamegpu/runtime/HostAgentAPI.cu ${FLAMEGPU_ROOT}/src/flamegpu/runtime/messaging/BruteForce.cu diff --git a/src/flamegpu/gpu/CUDAAgent.cu b/src/flamegpu/gpu/CUDAAgent.cu index 2b88ab75d..0a0c2b30e 100644 --- a/src/flamegpu/gpu/CUDAAgent.cu +++ b/src/flamegpu/gpu/CUDAAgent.cu @@ -34,8 +34,8 @@ using std::experimental::filesystem::v1::path; #include "flamegpu/model/AgentDescription.h" #include "flamegpu/model/AgentFunctionDescription.h" -#include "flamegpu/runtime/cuRVE/curve.h" -#include "flamegpu/runtime/cuRVE/curve_rtc.h" +#include "flamegpu/runtime/detail/curve/curve.h" +#include "flamegpu/runtime/detail/curve/curve_rtc.h" #include "flamegpu/gpu/CUDAScatter.h" #include "flamegpu/util/detail/compute_capability.cuh" #include "flamegpu/util/nvtx.h" @@ -94,9 +94,9 @@ void CUDAAgent::mapRuntimeVariables(const AgentFunctionData& func, const unsigne agent_description.name.c_str(), func.initial_state.c_str()); } - const Curve::VariableHash agent_hash = Curve::variableRuntimeHash(agent_description.name.c_str()); - const Curve::VariableHash func_hash = Curve::variableRuntimeHash(func.name.c_str()); - auto &curve = Curve::getInstance(); + const detail::curve::Curve::VariableHash agent_hash = detail::curve::Curve::variableRuntimeHash(agent_description.name.c_str()); + const detail::curve::Curve::VariableHash func_hash = detail::curve::Curve::variableRuntimeHash(func.name.c_str()); + auto &curve = detail::curve::Curve::getInstance(); const unsigned int agent_count = this->getStateSize(func.initial_state); // loop through the agents variables to map each variable name using cuRVE for (const auto &mmp : agent_description.variables) { @@ -104,7 +104,7 @@ void CUDAAgent::mapRuntimeVariables(const AgentFunctionData& func, const unsigne void* d_ptr = sm->second->getVariablePointer(mmp.first); // map using curve - const Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str()); + const detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::variableRuntimeHash(mmp.first.c_str()); // get the agent variable size const size_t type_size = mmp.second.type_size * mmp.second.elements; @@ -112,9 +112,9 @@ void CUDAAgent::mapRuntimeVariables(const AgentFunctionData& func, const unsigne // maximum population num if (func.func || func.condition) { #ifdef _DEBUG - const Curve::Variable cv = curve.registerVariableByHash(var_hash + agent_hash + func_hash + instance_id, d_ptr, type_size, agent_count); - if (cv != static_cast((var_hash + agent_hash + func_hash + instance_id)%Curve::MAX_VARIABLES)) { - fprintf(stderr, "Curve Warning: Agent Function '%s' Variable '%s' has a collision and may work improperly.\n", func.name.c_str(), mmp.first.c_str()); + const detail::curve::Curve::Variable cv = curve.registerVariableByHash(var_hash + agent_hash + func_hash + instance_id, d_ptr, type_size, agent_count); + if (cv != static_cast((var_hash + agent_hash + func_hash + instance_id)%detail::curve::Curve::MAX_VARIABLES)) { + fprintf(stderr, "detail::curve::Curve Warning: Agent Function '%s' Variable '%s' has a collision and may work improperly.\n", func.name.c_str(), mmp.first.c_str()); } #else curve.registerVariableByHash(var_hash + agent_hash + func_hash + instance_id, d_ptr, type_size, agent_count); @@ -150,16 +150,16 @@ void CUDAAgent::unmapRuntimeVariables(const AgentFunctionData& func, const unsig agent_description.name.c_str(), func.initial_state.c_str()); } - const Curve::VariableHash agent_hash = Curve::variableRuntimeHash(agent_description.name.c_str()); - const Curve::VariableHash func_hash = Curve::variableRuntimeHash(func.name.c_str()); + const detail::curve::Curve::VariableHash agent_hash = detail::curve::Curve::variableRuntimeHash(agent_description.name.c_str()); + const detail::curve::Curve::VariableHash func_hash = detail::curve::Curve::variableRuntimeHash(func.name.c_str()); // loop through the agents variables to map each variable name using cuRVE for (const auto &mmp : agent_description.variables) { // get a device pointer for the agent variable name // void* d_ptr = sm->second->getAgentListVariablePointer(mmp.first); // unmap using curve - const Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str()); - Curve::getInstance().unregisterVariableByHash(var_hash + agent_hash + func_hash + instance_id); + const detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::variableRuntimeHash(mmp.first.c_str()); + detail::curve::Curve::getInstance().unregisterVariableByHash(var_hash + agent_hash + func_hash + instance_id); } // No current need to unmap RTC variables as they are specific to the agent functions and thus do not persist beyond the scope of a single function @@ -413,13 +413,13 @@ void CUDAAgent::mapNewRuntimeVariables(const CUDAAgent& func_agent, const AgentF maxLen, 0); // Map variables to curve - const Curve::VariableHash _agent_birth_hash = Curve::variableRuntimeHash("_agent_birth"); - const Curve::VariableHash func_hash = Curve::variableRuntimeHash(func.name.c_str()); - auto &curve = Curve::getInstance(); + const detail::curve::Curve::VariableHash _agent_birth_hash = detail::curve::Curve::variableRuntimeHash("_agent_birth"); + const detail::curve::Curve::VariableHash func_hash = detail::curve::Curve::variableRuntimeHash(func.name.c_str()); + auto &curve = detail::curve::Curve::getInstance(); // loop through the agents variables to map each variable name using cuRVE for (const auto &mmp : agent_description.variables) { // map using curve - const Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str()); + const detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::variableRuntimeHash(mmp.first.c_str()); // get the agent variable size const size_t type_size = mmp.second.type_size * mmp.second.elements; @@ -438,9 +438,9 @@ void CUDAAgent::mapNewRuntimeVariables(const CUDAAgent& func_agent, const AgentF // maximum population num if (func.func) { #ifdef _DEBUG - const Curve::Variable cv = curve.registerVariableByHash(var_hash + (_agent_birth_hash ^ func_hash) + instance_id, d_ptr, type_size, maxLen); - if (cv != static_cast((var_hash + (_agent_birth_hash ^ func_hash) + instance_id)%Curve::MAX_VARIABLES)) { - fprintf(stderr, "Curve Warning: Agent Function '%s' New Agent Variable '%s' has a collision and may work improperly.\n", func.name.c_str(), mmp.first.c_str()); + const detail::curve::Curve::Variable cv = curve.registerVariableByHash(var_hash + (_agent_birth_hash ^ func_hash) + instance_id, d_ptr, type_size, maxLen); + if (cv != static_cast((var_hash + (_agent_birth_hash ^ func_hash) + instance_id)%detail::curve::Curve::MAX_VARIABLES)) { + fprintf(stderr, "detail::curve::Curve Warning: Agent Function '%s' New Agent Variable '%s' has a collision and may work improperly.\n", func.name.c_str(), mmp.first.c_str()); } #else curve.registerVariableByHash(var_hash + (_agent_birth_hash ^ func_hash) + instance_id, d_ptr, type_size, maxLen); @@ -472,13 +472,13 @@ void CUDAAgent::unmapNewRuntimeVariables(const AgentFunctionData& func, const un if (!func.func) return; // Unmap curve - const Curve::VariableHash _agent_birth_hash = Curve::variableRuntimeHash("_agent_birth"); - const Curve::VariableHash func_hash = Curve::variableRuntimeHash(func.name.c_str()); - auto &curve = Curve::getInstance(); + const detail::curve::Curve::VariableHash _agent_birth_hash = detail::curve::Curve::variableRuntimeHash("_agent_birth"); + const detail::curve::Curve::VariableHash func_hash = detail::curve::Curve::variableRuntimeHash(func.name.c_str()); + auto &curve = detail::curve::Curve::getInstance(); // loop through the agents variables to map each variable name using cuRVE for (const auto &mmp : agent_description.variables) { // unmap using curve - const Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str()); + const detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::variableRuntimeHash(mmp.first.c_str()); curve.unregisterVariableByHash(var_hash + (_agent_birth_hash ^ func_hash) + instance_id); // no need to unmap RTC variables @@ -519,7 +519,7 @@ void CUDAAgent::clearFunctionCondition(const std::string &state) { void CUDAAgent::addInstantitateRTCFunction(const AgentFunctionData& func, bool function_condition) { // Generate the dynamic curve header - CurveRTCHost &curve_header = *rtc_header_map.emplace(function_condition ? func.name + "_condition" : func.name, std::make_unique()).first->second; + detail::curve::CurveRTCHost &curve_header = *rtc_header_map.emplace(function_condition ? func.name + "_condition" : func.name, std::make_unique()).first->second; // set agent function variables in rtc curve for (const auto& mmp : func.parent.lock()->variables) { @@ -641,7 +641,7 @@ const jitify::experimental::KernelInstantiation& CUDAAgent::getRTCInstantiation( return *mm->second; } -CurveRTCHost& CUDAAgent::getRTCHeader(const std::string& function_name) const { +detail::curve::CurveRTCHost& CUDAAgent::getRTCHeader(const std::string& function_name) const { CUDARTCHeaderMap::const_iterator mm = rtc_header_map.find(function_name); if (mm == rtc_header_map.end()) { THROW exception::InvalidAgentFunc("Function name '%s' is not a runtime compiled agent function in agent '%s', " diff --git a/src/flamegpu/gpu/CUDAMessage.cu b/src/flamegpu/gpu/CUDAMessage.cu index a715cbbae..cca8cad44 100644 --- a/src/flamegpu/gpu/CUDAMessage.cu +++ b/src/flamegpu/gpu/CUDAMessage.cu @@ -9,7 +9,7 @@ #include "flamegpu/runtime/messaging/BruteForce.h" #include "flamegpu/model/AgentFunctionDescription.h" -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #include "flamegpu/model/AgentDescription.h" #include "flamegpu/runtime/messaging.h" @@ -98,17 +98,17 @@ void CUDAMessage::mapReadRuntimeVariables(const AgentFunctionData& func, const C const std::string message_name = message_description.name; - const Curve::VariableHash message_hash = Curve::getInstance().variableRuntimeHash(message_name.c_str()); - const Curve::VariableHash agent_hash = Curve::getInstance().variableRuntimeHash(func.parent.lock()->name.c_str()); - const Curve::VariableHash func_hash = Curve::getInstance().variableRuntimeHash(func.name.c_str()); - auto &curve = Curve::getInstance(); + const detail::curve::Curve::VariableHash message_hash = detail::curve::Curve::getInstance().variableRuntimeHash(message_name.c_str()); + const detail::curve::Curve::VariableHash agent_hash = detail::curve::Curve::getInstance().variableRuntimeHash(func.parent.lock()->name.c_str()); + const detail::curve::Curve::VariableHash func_hash = detail::curve::Curve::getInstance().variableRuntimeHash(func.name.c_str()); + auto &curve = detail::curve::Curve::getInstance(); // loop through the message variables to map each variable name using cuRVE for (const auto &mmp : message_description.variables) { // get a device pointer for the message variable name void* d_ptr = message_list->getReadMessageListVariablePointer(mmp.first); // map using curve - Curve::VariableHash var_hash = Curve::getInstance().variableRuntimeHash(mmp.first.c_str()); + detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::getInstance().variableRuntimeHash(mmp.first.c_str()); // get the message variable size size_t size = mmp.second.type_size; @@ -117,9 +117,9 @@ void CUDAMessage::mapReadRuntimeVariables(const AgentFunctionData& func, const C // maximum population size unsigned int length = this->getMessageCount(); // check to see if it is equal to pop #ifdef _DEBUG - const Curve::Variable cv = curve.registerVariableByHash(var_hash + agent_hash + func_hash + message_hash + instance_id, d_ptr, size, length); - if (cv != static_cast((var_hash + agent_hash + func_hash + message_hash + instance_id)%Curve::MAX_VARIABLES)) { - fprintf(stderr, "Curve Warning: Agent Function '%s' Message In Variable '%s' has a collision and may work improperly.\n", message_name.c_str(), mmp.first.c_str()); + const detail::curve::Curve::Variable cv = curve.registerVariableByHash(var_hash + agent_hash + func_hash + message_hash + instance_id, d_ptr, size, length); + if (cv != static_cast((var_hash + agent_hash + func_hash + message_hash + instance_id)%detail::curve::Curve::MAX_VARIABLES)) { + fprintf(stderr, "detail::curve::Curve Warning: Agent Function '%s' Message In Variable '%s' has a collision and may work improperly.\n", message_name.c_str(), mmp.first.c_str()); } #else curve.registerVariableByHash(var_hash + agent_hash + func_hash + message_hash + instance_id, d_ptr, size, length); @@ -149,17 +149,17 @@ void CUDAMessage::mapWriteRuntimeVariables(const AgentFunctionData& func, const const std::string message_name = message_description.name; - const Curve::VariableHash message_hash = Curve::getInstance().variableRuntimeHash(message_name.c_str()); - const Curve::VariableHash agent_hash = Curve::getInstance().variableRuntimeHash(func.parent.lock()->name.c_str()); - const Curve::VariableHash func_hash = Curve::getInstance().variableRuntimeHash(func.name.c_str()); - auto &curve = Curve::getInstance(); + const detail::curve::Curve::VariableHash message_hash = detail::curve::Curve::getInstance().variableRuntimeHash(message_name.c_str()); + const detail::curve::Curve::VariableHash agent_hash = detail::curve::Curve::getInstance().variableRuntimeHash(func.parent.lock()->name.c_str()); + const detail::curve::Curve::VariableHash func_hash = detail::curve::Curve::getInstance().variableRuntimeHash(func.name.c_str()); + auto &curve = detail::curve::Curve::getInstance(); // loop through the message variables to map each variable name using cuRVE for (const auto &mmp : message_description.variables) { // get a device pointer for the message variable name void* d_ptr = message_list->getWriteMessageListVariablePointer(mmp.first); // map using curve - Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str()); + detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::variableRuntimeHash(mmp.first.c_str()); // get the message variable size size_t size = mmp.second.type_size; @@ -168,9 +168,9 @@ void CUDAMessage::mapWriteRuntimeVariables(const AgentFunctionData& func, const // maximum population size unsigned int length = writeLen; // check to see if it is equal to pop #ifdef _DEBUG - const Curve::Variable cv = curve.registerVariableByHash(var_hash + agent_hash + func_hash + message_hash + instance_id, d_ptr, size, length); - if (cv != static_cast((var_hash + agent_hash + func_hash + message_hash + instance_id)%Curve::MAX_VARIABLES)) { - fprintf(stderr, "Curve Warning: Agent Function '%s' Message '%s' Out? Variable '%s' has a collision and may work improperly.\n", func.name.c_str(), message_name.c_str(), mmp.first.c_str()); + const detail::curve::Curve::Variable cv = curve.registerVariableByHash(var_hash + agent_hash + func_hash + message_hash + instance_id, d_ptr, size, length); + if (cv != static_cast((var_hash + agent_hash + func_hash + message_hash + instance_id)%detail::curve::Curve::MAX_VARIABLES)) { + fprintf(stderr, "detail::curve::Curve Warning: Agent Function '%s' Message '%s' Out? Variable '%s' has a collision and may work improperly.\n", func.name.c_str(), message_name.c_str(), mmp.first.c_str()); } #else curve.registerVariableByHash(var_hash + agent_hash + func_hash + message_hash + instance_id, d_ptr, size, length); @@ -198,14 +198,14 @@ void CUDAMessage::unmapRuntimeVariables(const AgentFunctionData& func, const uns } const std::string message_name = message_description.name; - const Curve::VariableHash message_hash = Curve::getInstance().variableRuntimeHash(message_name.c_str()); - const Curve::VariableHash agent_hash = Curve::getInstance().variableRuntimeHash(func.parent.lock()->name.c_str()); - const Curve::VariableHash func_hash = Curve::getInstance().variableRuntimeHash(func.name.c_str()); - auto &curve = Curve::getInstance(); + const detail::curve::Curve::VariableHash message_hash = detail::curve::Curve::getInstance().variableRuntimeHash(message_name.c_str()); + const detail::curve::Curve::VariableHash agent_hash = detail::curve::Curve::getInstance().variableRuntimeHash(func.parent.lock()->name.c_str()); + const detail::curve::Curve::VariableHash func_hash = detail::curve::Curve::getInstance().variableRuntimeHash(func.name.c_str()); + auto &curve = detail::curve::Curve::getInstance(); // loop through the message variables to map each variable name using cuRVE for (const auto &mmp : message_description.variables) { // unmap using curve - Curve::VariableHash var_hash = Curve::variableRuntimeHash(mmp.first.c_str()); + detail::curve::Curve::VariableHash var_hash = detail::curve::Curve::variableRuntimeHash(mmp.first.c_str()); curve.unregisterVariableByHash(var_hash + agent_hash + func_hash + message_hash + instance_id); } } diff --git a/src/flamegpu/gpu/CUDASimulation.cu b/src/flamegpu/gpu/CUDASimulation.cu index d18e2bdcf..a8b452ba1 100644 --- a/src/flamegpu/gpu/CUDASimulation.cu +++ b/src/flamegpu/gpu/CUDASimulation.cu @@ -16,7 +16,7 @@ #include "flamegpu/util/detail/compute_capability.cuh" #include "flamegpu/util/detail/SignalHandlers.h" #include "flamegpu/util/detail/CUDAEventTimer.cuh" -#include "flamegpu/runtime/cuRVE/curve_rtc.h" +#include "flamegpu/runtime/detail/curve/curve_rtc.h" #include "flamegpu/runtime/HostFunctionCallback.h" #include "flamegpu/gpu/CUDAAgent.h" #include "flamegpu/gpu/CUDAMessage.h" @@ -200,7 +200,7 @@ CUDASimulation::~CUDASimulation() { // Could mutex it with init simulation cuda stuff, but really seems unlikely gpuErrchk(cudaDeviceReset()); EnvironmentManager::getInstance().purge(); - Curve::getInstance().purge(); + detail::curve::Curve::getInstance().purge(); } } if (t_device_id != deviceInitialised) { @@ -421,9 +421,9 @@ void CUDASimulation::stepLayer(const std::shared_ptr& layer, const un int gridSize = 0; // The actual grid size needed, based on input size // Agent function condition kernel wrapper args - Curve::NamespaceHash agentname_hash = Curve::variableRuntimeHash(agent_name.c_str()); - Curve::NamespaceHash funcname_hash = Curve::variableRuntimeHash(func_name.c_str()); - Curve::NamespaceHash agent_func_name_hash = agentname_hash + funcname_hash + instance_id; + detail::curve::Curve::NamespaceHash agentname_hash = detail::curve::Curve::variableRuntimeHash(agent_name.c_str()); + detail::curve::Curve::NamespaceHash funcname_hash = detail::curve::Curve::variableRuntimeHash(func_name.c_str()); + detail::curve::Curve::NamespaceHash agent_func_name_hash = agentname_hash + funcname_hash + instance_id; curandState *t_rng = d_rng + totalThreads; unsigned int *scanFlag_agentDeath = this->singletons->scatter.Scan().Config(CUDAScanCompaction::Type::AGENT_DEATH, streamIdx).d_ptrs.scan_flag; unsigned int sm_size = 0; @@ -625,12 +625,12 @@ void CUDASimulation::stepLayer(const std::shared_ptr& layer, const un id_t *d_agentOut_nextID = nullptr; std::string agent_name = func_agent->name; std::string func_name = func_des->name; - Curve::NamespaceHash agentname_hash = Curve::variableRuntimeHash(agent_name.c_str()); - Curve::NamespaceHash funcname_hash = Curve::variableRuntimeHash(func_name.c_str()); - Curve::NamespaceHash agent_func_name_hash = agentname_hash + funcname_hash + instance_id; - Curve::NamespaceHash message_name_inp_hash = 0; - Curve::NamespaceHash message_name_outp_hash = 0; - Curve::NamespaceHash agentoutput_hash = 0; + detail::curve::Curve::NamespaceHash agentname_hash = detail::curve::Curve::variableRuntimeHash(agent_name.c_str()); + detail::curve::Curve::NamespaceHash funcname_hash = detail::curve::Curve::variableRuntimeHash(func_name.c_str()); + detail::curve::Curve::NamespaceHash agent_func_name_hash = agentname_hash + funcname_hash + instance_id; + detail::curve::Curve::NamespaceHash message_name_inp_hash = 0; + detail::curve::Curve::NamespaceHash message_name_outp_hash = 0; + detail::curve::Curve::NamespaceHash agentoutput_hash = 0; // check if a function has an input message if (auto im = func_des->message_input.lock()) { @@ -638,7 +638,7 @@ void CUDASimulation::stepLayer(const std::shared_ptr& layer, const un const CUDAMessage& cuda_message = getCUDAMessage(inpMessage_name); // hash message name - message_name_inp_hash = Curve::variableRuntimeHash(inpMessage_name.c_str()); + message_name_inp_hash = detail::curve::Curve::variableRuntimeHash(inpMessage_name.c_str()); d_in_messagelist_metadata = cuda_message.getMetaDataDevicePtr(); } @@ -649,13 +649,13 @@ void CUDASimulation::stepLayer(const std::shared_ptr& layer, const un const CUDAMessage& cuda_message = getCUDAMessage(outpMessage_name); // hash message name - message_name_outp_hash = Curve::variableRuntimeHash(outpMessage_name.c_str()); + message_name_outp_hash = detail::curve::Curve::variableRuntimeHash(outpMessage_name.c_str()); d_out_messagelist_metadata = cuda_message.getMetaDataDevicePtr(); } // check if a function has agent output if (auto oa = func_des->agent_output.lock()) { - agentoutput_hash = (Curve::variableRuntimeHash("_agent_birth") ^ funcname_hash) + instance_id; + agentoutput_hash = (detail::curve::Curve::variableRuntimeHash("_agent_birth") ^ funcname_hash) + instance_id; CUDAAgent& output_agent = getCUDAAgent(oa->name); d_agentOut_nextID = output_agent.getDeviceNextID(); } @@ -1275,7 +1275,7 @@ void CUDASimulation::initialiseSingletons() { gpuErrchk(cudaMemcpyFromSymbol(&DEVICE_HAS_RESET_CHECK, DEVICE_HAS_RESET, sizeof(unsigned int))); if (DEVICE_HAS_RESET_CHECK == DEVICE_HAS_RESET_FLAG) { // Device has been reset, purge host mirrors of static objects/singletons - Curve::getInstance().purge(); + detail::curve::Curve::getInstance().purge(); if (singletons) { singletons->rng.purge(); singletons->scatter.purge(); @@ -1289,7 +1289,7 @@ void CUDASimulation::initialiseSingletons() { maps_lock.unlock(); // Get references to all required singleton and store in the instance. singletons = new Singletons( - Curve::getInstance(), + detail::curve::Curve::getInstance(), EnvironmentManager::getInstance()); // Reinitialise random for this simulation instance diff --git a/src/flamegpu/model/AgentFunctionData.cpp b/src/flamegpu/model/AgentFunctionData.cpp index 7a0c25db0..ddac64fa0 100644 --- a/src/flamegpu/model/AgentFunctionData.cpp +++ b/src/flamegpu/model/AgentFunctionData.cpp @@ -2,7 +2,7 @@ #include "flamegpu/model/AgentDescription.h" #include "flamegpu/model/AgentFunctionDescription.h" -#include "flamegpu/runtime/cuRVE/curve_rtc.h" +#include "flamegpu/runtime/detail/curve/curve_rtc.h" #include "flamegpu/util/detail/cxxname.hpp" namespace flamegpu { @@ -64,22 +64,22 @@ AgentFunctionData::AgentFunctionData(const std::shared_ptr &mod if (_m != model->messages.end()) { message_input = _m->second; } - } else if (util::detail::cxxname::getUnqualifiedName(other.msg_in_type) != util::detail::cxxname::getUnqualifiedName(CurveRTCHost::demangle(std::type_index(typeid(MsgNone))))) { + } else if (util::detail::cxxname::getUnqualifiedName(other.msg_in_type) != util::detail::cxxname::getUnqualifiedName(detail::curve::CurveRTCHost::demangle(std::type_index(typeid(MsgNone))))) { THROW exception::InvalidMessageType( "Function '%s' is missing bound input message of type '%s', type provided was '%s'.", other.name.c_str(), util::detail::cxxname::getUnqualifiedName(other.msg_in_type).c_str(), - util::detail::cxxname::getUnqualifiedName(CurveRTCHost::demangle(std::type_index(typeid(MsgNone)))).c_str()); + util::detail::cxxname::getUnqualifiedName(detail::curve::CurveRTCHost::demangle(std::type_index(typeid(MsgNone)))).c_str()); } if (auto a = other.message_output.lock()) { auto _m = model->messages.find(a->name); if (_m != model->messages.end()) { message_output = _m->second; } - } else if (util::detail::cxxname::getUnqualifiedName(other.msg_out_type) != util::detail::cxxname::getUnqualifiedName(CurveRTCHost::demangle(std::type_index(typeid(MsgNone))))) { + } else if (util::detail::cxxname::getUnqualifiedName(other.msg_out_type) != util::detail::cxxname::getUnqualifiedName(detail::curve::CurveRTCHost::demangle(std::type_index(typeid(MsgNone))))) { THROW exception::InvalidMessageType( "Function '%s' is missing bound output message of type '%s', type provided was '%s'.", other.name.c_str(), util::detail::cxxname::getUnqualifiedName(other.msg_out_type).c_str(), - util::detail::cxxname::getUnqualifiedName(CurveRTCHost::demangle(std::type_index(typeid(MsgNone)))).c_str()); + util::detail::cxxname::getUnqualifiedName(detail::curve::CurveRTCHost::demangle(std::type_index(typeid(MsgNone)))).c_str()); } if (auto a = other.agent_output.lock()) { auto _a = model->agents.find(a->name); diff --git a/src/flamegpu/model/AgentFunctionDescription.cpp b/src/flamegpu/model/AgentFunctionDescription.cpp index ae470e325..06e914c12 100644 --- a/src/flamegpu/model/AgentFunctionDescription.cpp +++ b/src/flamegpu/model/AgentFunctionDescription.cpp @@ -141,7 +141,7 @@ void AgentFunctionDescription::setMessageInput(const std::string &message_name) if (a != mdl->messages.end()) { // Just compare the classname is the same, to allow for the various approaches to namespace use. This should only be required for RTC functions. auto msg_in_classname = util::detail::cxxname::getUnqualifiedName(this->function->msg_in_type); - auto demangledClassName = util::detail::cxxname::getUnqualifiedName(CurveRTCHost::demangle(a->second->getType())); + auto demangledClassName = util::detail::cxxname::getUnqualifiedName(detail::curve::CurveRTCHost::demangle(a->second->getType())); if (msg_in_classname == demangledClassName) { this->function->message_input = a->second; } else { @@ -177,7 +177,7 @@ void AgentFunctionDescription::setMessageInput(MsgBruteForce::Description &messa if (a->second->description.get() == &message) { // Just compare the classname is the same, to allow for the various approaches to namespace use. This should only be required for RTC functions. auto msg_in_classname = util::detail::cxxname::getUnqualifiedName(this->function->msg_in_type); - auto demangledClassName = util::detail::cxxname::getUnqualifiedName(CurveRTCHost::demangle(a->second->getType())); + auto demangledClassName = util::detail::cxxname::getUnqualifiedName(detail::curve::CurveRTCHost::demangle(a->second->getType())); if (msg_in_classname == demangledClassName) { this->function->message_input = a->second; } else { @@ -219,7 +219,7 @@ void AgentFunctionDescription::setMessageOutput(const std::string &message_name) if (a != mdl->messages.end()) { // Just compare the classname is the same, to allow for the various approaches to namespace use. This should only be required for RTC functions. auto msg_out_classname = util::detail::cxxname::getUnqualifiedName(this->function->msg_out_type); - auto demangledClassName = util::detail::cxxname::getUnqualifiedName(CurveRTCHost::demangle(a->second->getType())); + auto demangledClassName = util::detail::cxxname::getUnqualifiedName(detail::curve::CurveRTCHost::demangle(a->second->getType())); if (msg_out_classname == demangledClassName) { this->function->message_output = a->second; if (this->function->message_output_optional) { @@ -264,7 +264,7 @@ void AgentFunctionDescription::setMessageOutput(MsgBruteForce::Description &mess if (a->second->description.get() == &message) { // Just compare the classname is the same, to allow for the various approaches to namespace use. This should only be required for RTC functions. auto msg_out_classname = util::detail::cxxname::getUnqualifiedName(this->function->msg_out_type); - auto demangledClassName = util::detail::cxxname::getUnqualifiedName(CurveRTCHost::demangle(a->second->getType())); + auto demangledClassName = util::detail::cxxname::getUnqualifiedName(detail::curve::CurveRTCHost::demangle(a->second->getType())); if (msg_out_classname == demangledClassName) { this->function->message_output = a->second; if (this->function->message_output_optional) { diff --git a/src/flamegpu/runtime/cuRVE/curve.cu b/src/flamegpu/runtime/detail/curve/curve.cu similarity index 98% rename from src/flamegpu/runtime/cuRVE/curve.cu rename to src/flamegpu/runtime/detail/curve/curve.cu index d47d00f94..b2af8c050 100644 --- a/src/flamegpu/runtime/cuRVE/curve.cu +++ b/src/flamegpu/runtime/detail/curve/curve.cu @@ -5,13 +5,14 @@ #include #include -#include "flamegpu/runtime/cuRVE/curve.h" +#include "flamegpu/runtime/detail/curve/curve.h" #include "flamegpu/gpu/detail/CUDAErrorChecking.cuh" #include "flamegpu/util/nvtx.h" namespace flamegpu { +namespace detail { namespace curve { namespace detail { /** @@ -33,7 +34,6 @@ namespace detail { */ __constant__ unsigned int d_lengths[Curve::MAX_VARIABLES]; } // namespace detail -} // namespace curve std::mutex Curve::instance_mutex; @@ -210,4 +210,6 @@ Curve& Curve::getInstance() { return *(instances.emplace(device_id, std::unique_ptr(new Curve())).first->second); } +} // namespace curve +} // namespace detail } // namespace flamegpu diff --git a/src/flamegpu/runtime/cuRVE/curve_rtc.cpp b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp similarity index 94% rename from src/flamegpu/runtime/cuRVE/curve_rtc.cpp rename to src/flamegpu/runtime/detail/curve/curve_rtc.cpp index cf0cd3a19..c887dd5f2 100644 --- a/src/flamegpu/runtime/cuRVE/curve_rtc.cpp +++ b/src/flamegpu/runtime/detail/curve/curve_rtc.cpp @@ -1,6 +1,6 @@ #include -#include "flamegpu/runtime/cuRVE/curve_rtc.h" +#include "flamegpu/runtime/detail/curve/curve_rtc.h" #include "flamegpu/exception/FLAMEGPUException.h" #include "flamegpu/runtime/utility/EnvironmentManager.cuh" @@ -14,6 +14,8 @@ #endif namespace flamegpu { +namespace detail { +namespace curve { const char* CurveRTCHost::curve_rtc_dynamic_h_template = R"###(dynamic/curve_rtc_dynamic.h @@ -46,6 +48,9 @@ __device__ bool strings_equal(const char(&a)[N], const char(&b)[M]) { return false; } +namespace detail { +namespace curve { + /** * Dynamically generated version of Curve without hashing * Both environment data, and curve variable ptrs are stored in this buffer @@ -158,6 +163,8 @@ __device__ __forceinline__ void Curve::setNewAgentArrayVariable(const char(&name $DYNAMIC_SETNEWAGENTARRAYVARIABLE_IMPL } +} // namespace curve +} // namespace detail } // namespace flamegpu // has to be included after definition of curve namespace @@ -348,7 +355,7 @@ void CurveRTCHost::initHeaderEnvironment() { getEnvVariableImpl << " return 0;\n"; getEnvVariableImpl << " }\n"; getEnvVariableImpl << "#endif\n"; - getEnvVariableImpl << " return *reinterpret_cast(reinterpret_cast(" << getVariableSymbolName() <<" + " << props.offset << "));\n"; + getEnvVariableImpl << " return *reinterpret_cast(reinterpret_cast(flamegpu::detail::curve::" << getVariableSymbolName() <<" + " << props.offset << "));\n"; getEnvVariableImpl << " };\n"; } } @@ -378,7 +385,7 @@ void CurveRTCHost::initHeaderEnvironment() { getEnvArrayVariableImpl << " return 0;\n"; getEnvArrayVariableImpl << " }\n"; getEnvArrayVariableImpl << "#endif\n"; - getEnvArrayVariableImpl << " return reinterpret_cast(reinterpret_cast(" << getVariableSymbolName() <<" + " << props.offset << "))[index];\n"; + getEnvArrayVariableImpl << " return reinterpret_cast(reinterpret_cast(flamegpu::detail::curve::" << getVariableSymbolName() <<" + " << props.offset << "))[index];\n"; getEnvArrayVariableImpl << " };\n"; } } @@ -417,7 +424,7 @@ void CurveRTCHost::initHeaderSetters() { setAgentVariableImpl << " return;\n"; setAgentVariableImpl << " }\n"; setAgentVariableImpl << "#endif\n"; - setAgentVariableImpl << " (*static_cast(static_cast(" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << ")))[index] = (T) variable;\n"; + setAgentVariableImpl << " (*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << ")))[index] = (T) variable;\n"; setAgentVariableImpl << " return;\n"; setAgentVariableImpl << " }\n"; } else { ++ct; } @@ -441,7 +448,7 @@ void CurveRTCHost::initHeaderSetters() { setMessageVariableImpl << " return;\n"; setMessageVariableImpl << " }\n"; setMessageVariableImpl << "#endif\n"; - setMessageVariableImpl << " (*static_cast(static_cast(" << getVariableSymbolName() << " + " << msgOut_data_offset + (ct++ * sizeof(void*)) << ")))[index] = (T) variable;\n"; + setMessageVariableImpl << " (*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << msgOut_data_offset + (ct++ * sizeof(void*)) << ")))[index] = (T) variable;\n"; setMessageVariableImpl << " return;\n"; setMessageVariableImpl << " }\n"; } else { ++ct; } @@ -465,7 +472,7 @@ void CurveRTCHost::initHeaderSetters() { setNewAgentVariableImpl << " return;\n"; setNewAgentVariableImpl << " }\n"; setNewAgentVariableImpl << "#endif\n"; - setNewAgentVariableImpl << " (*static_cast(static_cast(" << getVariableSymbolName() << " + " << newAgent_data_offset + (ct++ * sizeof(void*)) << ")))[index] = (T) variable;\n"; + setNewAgentVariableImpl << " (*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << newAgent_data_offset + (ct++ * sizeof(void*)) << ")))[index] = (T) variable;\n"; setNewAgentVariableImpl << " return;\n"; setNewAgentVariableImpl << " }\n"; } else { ++ct; } @@ -497,7 +504,7 @@ void CurveRTCHost::initHeaderSetters() { setAgentArrayVariableImpl << " return;\n"; setAgentArrayVariableImpl << " }\n"; setAgentArrayVariableImpl << "#endif\n"; - setAgentArrayVariableImpl << " (*static_cast(static_cast(" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << ")))[i] = (T) variable;\n"; + setAgentArrayVariableImpl << " (*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << ")))[i] = (T) variable;\n"; setAgentArrayVariableImpl << " return;\n"; setAgentArrayVariableImpl << " }\n"; } else { ++ct; } @@ -529,7 +536,7 @@ void CurveRTCHost::initHeaderSetters() { setNewAgentArrayVariableImpl << " return;\n"; setNewAgentArrayVariableImpl << " }\n"; setNewAgentArrayVariableImpl << "#endif\n"; - setNewAgentArrayVariableImpl << " (*static_cast(static_cast(" << getVariableSymbolName() << " + " << newAgent_data_offset + (ct++ * sizeof(void*)) << ")))[i] = (T) variable;\n"; + setNewAgentArrayVariableImpl << " (*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << newAgent_data_offset + (ct++ * sizeof(void*)) << ")))[i] = (T) variable;\n"; setNewAgentArrayVariableImpl << " return;\n"; setNewAgentArrayVariableImpl << " }\n"; } else { ++ct; } @@ -555,7 +562,7 @@ void CurveRTCHost::initHeaderGetters() { getAgentVariableImpl << " return 0;\n"; getAgentVariableImpl << " }\n"; getAgentVariableImpl << "#endif\n"; - getAgentVariableImpl << " return (*static_cast(static_cast(" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << ")))[index];\n"; + getAgentVariableImpl << " return (*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << ")))[index];\n"; getAgentVariableImpl << " }\n"; } else { ++ct; } } @@ -579,7 +586,7 @@ void CurveRTCHost::initHeaderGetters() { getMessageVariableImpl << " return 0;\n"; getMessageVariableImpl << " }\n"; getMessageVariableImpl << "#endif\n"; - getMessageVariableImpl << " return (*static_cast(static_cast(" << getVariableSymbolName() << " + " << msgIn_data_offset + (ct++ * sizeof(void*)) << ")))[index];\n"; + getMessageVariableImpl << " return (*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << msgIn_data_offset + (ct++ * sizeof(void*)) << ")))[index];\n"; getMessageVariableImpl << " }\n"; } else { ++ct; } } @@ -603,7 +610,7 @@ void CurveRTCHost::initHeaderGetters() { getAgentVariableLDGImpl << " return 0;\n"; getAgentVariableLDGImpl << " }\n"; getAgentVariableLDGImpl << "#endif\n"; - getAgentVariableLDGImpl << " return (T) __ldg((*static_cast(static_cast(" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << "))) + index);\n"; + getAgentVariableLDGImpl << " return (T) __ldg((*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << "))) + index);\n"; getAgentVariableLDGImpl << " }\n"; } else { ++ct; } } @@ -627,7 +634,7 @@ void CurveRTCHost::initHeaderGetters() { getMessageVariableLDGImpl << " return 0;\n"; getMessageVariableLDGImpl << " }\n"; getMessageVariableLDGImpl << "#endif\n"; - getMessageVariableLDGImpl << " return (T) __ldg((*static_cast(static_cast(" << getVariableSymbolName() << " + " << msgIn_data_offset + (ct++ * sizeof(void*)) << "))) + index);\n"; + getMessageVariableLDGImpl << " return (T) __ldg((*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << msgIn_data_offset + (ct++ * sizeof(void*)) << "))) + index);\n"; getMessageVariableLDGImpl << " }\n"; } else { ++ct; } } @@ -659,7 +666,7 @@ void CurveRTCHost::initHeaderGetters() { getAgentArrayVariableImpl << " return 0;\n"; getAgentArrayVariableImpl << " }\n"; getAgentArrayVariableImpl << "#endif\n"; - getAgentArrayVariableImpl << " return (*static_cast(static_cast(" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << ")))[i];\n"; + getAgentArrayVariableImpl << " return (*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << ")))[i];\n"; getAgentArrayVariableImpl << " };\n"; } else { ++ct; } } @@ -691,7 +698,7 @@ void CurveRTCHost::initHeaderGetters() { getAgentArrayVariableLDGImpl << " return 0;\n"; getAgentArrayVariableLDGImpl << " }\n"; getAgentArrayVariableLDGImpl << "#endif\n"; - getAgentArrayVariableLDGImpl << " return (T) __ldg((*static_cast(static_cast(" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << "))) + i);\n"; + getAgentArrayVariableLDGImpl << " return (T) __ldg((*static_cast(static_cast(flamegpu::detail::curve::" << getVariableSymbolName() << " + " << agent_data_offset + (ct++ * sizeof(void*)) << "))) + i);\n"; getAgentArrayVariableLDGImpl << " };\n"; } else { ++ct; } } @@ -799,9 +806,11 @@ void CurveRTCHost::updateEnvCache(const char *env_ptr) { } void CurveRTCHost::updateDevice(const jitify::experimental::KernelInstantiation& instance) { // The namespace is required here, but not in other uses of getVariableSymbolName. - std::string cache_var_name = std::string("flamegpu::") + getVariableSymbolName(); + std::string cache_var_name = std::string("flamegpu::detail::curve::") + getVariableSymbolName(); CUdeviceptr d_var_ptr = instance.get_global_ptr(cache_var_name.c_str()); gpuErrchkDriverAPI(cuMemcpyHtoD(d_var_ptr, h_data_buffer, data_buffer_size)); } +} // namespace curve +} // namespace detail } // namespace flamegpu diff --git a/src/flamegpu/runtime/utility/EnvironmentManager.cu b/src/flamegpu/runtime/utility/EnvironmentManager.cu index 5daccf0b5..11aaeaa66 100644 --- a/src/flamegpu/runtime/utility/EnvironmentManager.cu +++ b/src/flamegpu/runtime/utility/EnvironmentManager.cu @@ -26,7 +26,7 @@ std::mutex EnvironmentManager::instance_mutex; const char EnvironmentManager::CURVE_NAMESPACE_STRING[23] = "ENVIRONMENT_PROPERTIES"; EnvironmentManager::EnvironmentManager() : - CURVE_NAMESPACE_HASH(Curve::variableRuntimeHash(CURVE_NAMESPACE_STRING)), + CURVE_NAMESPACE_HASH(detail::curve::Curve::variableRuntimeHash(CURVE_NAMESPACE_STRING)), nextFree(0), m_freeSpace(EnvironmentManager::MAX_BUFFER_SIZE), freeFragments(), @@ -87,7 +87,7 @@ void EnvironmentManager::init(const unsigned int &instance_id, const Environment "in EnvironmentManager::init()."); } // Defragment to rebuild it properly - defragment(Curve::getInstance(), &orderedProperties, {}, isPureRTC); + defragment(detail::curve::Curve::getInstance(), &orderedProperties, {}, isPureRTC); // Setup RTC version buildRTCOffsets(instance_id, instance_id, orderedProperties); } @@ -144,7 +144,7 @@ void EnvironmentManager::init(const unsigned int &instance_id, const Environment "in EnvironmentManager::init()."); } // Defragment to rebuild it properly - defragment(Curve::getInstance(), &orderedProperties, new_mapped_props, isPureRTC); + defragment(detail::curve::Curve::getInstance(), &orderedProperties, new_mapped_props, isPureRTC); // Setup RTC version buildRTCOffsets(instance_id, master_instance_id, orderedProperties); } @@ -171,13 +171,13 @@ void EnvironmentManager::initialiseDevice() { deviceInitialised = true; } } -void EnvironmentManager::free(Curve &curve, const unsigned int &instance_id) { +void EnvironmentManager::free(detail::curve::Curve &curve, const unsigned int &instance_id) { std::unique_lock lock(mutex); // Release regular properties for (auto &&i = properties.begin(); i != properties.end();) { if (i->first.first == instance_id) { // Release from CURVE - Curve::VariableHash cvh = toHash(i->first); + detail::curve::Curve::VariableHash cvh = toHash(i->first); curve.unregisterVariableByHash(cvh); // Drop from properties map i = properties.erase(i); @@ -189,7 +189,7 @@ void EnvironmentManager::free(Curve &curve, const unsigned int &instance_id) { for (auto &&i = mapped_properties.begin(); i != mapped_properties.end();) { if (i->first.first == instance_id) { // Release from CURVE - Curve::VariableHash cvh = toHash(i->first); + detail::curve::Curve::VariableHash cvh = toHash(i->first); curve.unregisterVariableByHash(cvh); // Drop from properties map i = mapped_properties.erase(i); @@ -227,8 +227,8 @@ EnvironmentManager::NamePair EnvironmentManager::toName(const unsigned int &inst /** * @note Not static, because eventually we might need to use curve singleton */ -Curve::VariableHash EnvironmentManager::toHash(const NamePair &name) const { - Curve::VariableHash var_cvh = Curve::variableRuntimeHash(name.second.c_str()); +detail::curve::Curve::VariableHash EnvironmentManager::toHash(const NamePair &name) const { + detail::curve::Curve::VariableHash var_cvh = detail::curve::Curve::variableRuntimeHash(name.second.c_str()); return CURVE_NAMESPACE_HASH + name.first + var_cvh; } @@ -276,7 +276,7 @@ void EnvironmentManager::newProperty(const NamePair &name, const char *ptr, cons } if (buffOffset == MAX_BUFFER_SIZE) { // buffOffset hasn't changed from init value // defragment() and retry using nextFree - defragment(Curve::getInstance()); + defragment(detail::curve::Curve::getInstance()); const ptrdiff_t alignmentOffset = nextFree % typeSize; const ptrdiff_t alignmentFix = alignmentOffset != 0 ? typeSize - alignmentOffset : 0; if (nextFree + alignmentFix + length < MAX_BUFFER_SIZE) { @@ -302,14 +302,14 @@ void EnvironmentManager::newProperty(const NamePair &name, const char *ptr, cons // Store data memcpy(hc_buffer + buffOffset, ptr, length); // Register in cuRVE - Curve::VariableHash cvh = toHash(name); - const auto CURVE_RESULT = Curve::getInstance().registerVariableByHash(cvh, reinterpret_cast(buffOffset), typeSize, elements); - if (CURVE_RESULT == Curve::UNKNOWN_VARIABLE) { + detail::curve::Curve::VariableHash cvh = toHash(name); + const auto CURVE_RESULT = detail::curve::Curve::getInstance().registerVariableByHash(cvh, reinterpret_cast(buffOffset), typeSize, elements); + if (CURVE_RESULT == detail::curve::Curve::UNKNOWN_VARIABLE) { THROW exception::CurveException("curveRegisterVariableByHash() returned UNKNOWN_CURVE_VARIABLE" "in EnvironmentManager::add()."); } #ifdef _DEBUG - if (CURVE_RESULT != static_cast(cvh%Curve::MAX_VARIABLES)) { + if (CURVE_RESULT != static_cast(cvh%detail::curve::Curve::MAX_VARIABLES)) { fprintf(stderr, "Curve Warning: Environment Property '%s' has a collision and may work improperly.\n", name.second.c_str()); } #endif @@ -318,9 +318,9 @@ void EnvironmentManager::newProperty(const NamePair &name, const char *ptr, cons } #ifdef _DEBUG -void EnvironmentManager::defragment(Curve &curve, const DefragMap * mergeProperties, std::set newmaps, bool isPureRTC) { +void EnvironmentManager::defragment(detail::curve::Curve &curve, const DefragMap * mergeProperties, std::set newmaps, bool isPureRTC) { #else -void EnvironmentManager::defragment(Curve & curve, const DefragMap * mergeProperties, std::set newmaps, bool) { +void EnvironmentManager::defragment(detail::curve::Curve & curve, const DefragMap * mergeProperties, std::set newmaps, bool) { #endif // Do not lock mutex here, do it in the calling method auto device_lock = std::unique_lock(device_mutex); @@ -360,7 +360,7 @@ void EnvironmentManager::defragment(Curve & curve, const DefragMap * mergeProper memcpy(t_buffer + buffOffset, i.data, i.length); t_properties.emplace(name, EnvProp(buffOffset, i.length, i.isConst, i.elements, i.type, i.rtc_offset)); // Update cuRVE (There isn't an update, so unregister and reregister) // TODO: curveGetVariableHandle()? - Curve::VariableHash cvh = toHash(name); + detail::curve::Curve::VariableHash cvh = toHash(name); // Only unregister variable if it's already registered if (!mergeProperties) { // Merge properties are only provided on 1st init, when vars can't be unregistered curve.unregisterVariableByHash(cvh); @@ -380,12 +380,12 @@ void EnvironmentManager::defragment(Curve & curve, const DefragMap * mergeProper } const auto CURVE_RESULT = curve.registerVariableByHash(cvh, reinterpret_cast(buffOffset), typeSize, i.elements); - if (CURVE_RESULT == Curve::UNKNOWN_VARIABLE) { + if (CURVE_RESULT == detail::curve::Curve::UNKNOWN_VARIABLE) { THROW exception::CurveException("curveRegisterVariableByHash() returned UNKNOWN_CURVE_VARIABLE, " "in EnvironmentManager::defragment()."); } #ifdef _DEBUG - if (!isPureRTC && CURVE_RESULT != static_cast(cvh%Curve::MAX_VARIABLES)) { + if (!isPureRTC && CURVE_RESULT != static_cast(cvh%detail::curve::Curve::MAX_VARIABLES)) { fprintf(stderr, "Curve Warning: Environment Property '%s' has a collision and may work improperly.\n", name.second.c_str()); } #endif @@ -408,7 +408,7 @@ void EnvironmentManager::defragment(Curve & curve, const DefragMap * mergeProper // Update cub for any mapped properties for (auto &mp : mapped_properties) { // Generate hash for the subproperty name - Curve::VariableHash cvh = toHash(mp.first); + detail::curve::Curve::VariableHash cvh = toHash(mp.first); // Unregister the property if it's already been registered if (newmaps.find(mp.first) == newmaps.end()) { curve.unregisterVariableByHash(cvh); @@ -418,12 +418,12 @@ void EnvironmentManager::defragment(Curve & curve, const DefragMap * mergeProper // Create the mapping const auto CURVE_RESULT = curve.registerVariableByHash(cvh, reinterpret_cast(masterprop.offset), masterprop.length / masterprop.elements, masterprop.elements); - if (CURVE_RESULT == Curve::UNKNOWN_VARIABLE) { + if (CURVE_RESULT == detail::curve::Curve::UNKNOWN_VARIABLE) { THROW exception::CurveException("curveRegisterVariableByHash() returned UNKNOWN_CURVE_VARIABLE, " "in EnvironmentManager::defragment()."); } #ifdef _DEBUG - if (!isPureRTC && CURVE_RESULT != static_cast(cvh%Curve::MAX_VARIABLES)) { + if (!isPureRTC && CURVE_RESULT != static_cast(cvh%detail::curve::Curve::MAX_VARIABLES)) { fprintf(stderr, "Curve Warning: Environment Property '%s' has a collision and may work improperly.\n", mp.first.second.c_str()); } #endif @@ -576,8 +576,8 @@ void EnvironmentManager::updateRTCValue(const NamePair &name) { void EnvironmentManager::removeProperty(const NamePair &name) { std::unique_lock lock(mutex); // Unregister in cuRVE - Curve::VariableHash cvh = toHash(name); - Curve::getInstance().unregisterVariableByHash(cvh); + detail::curve::Curve::VariableHash cvh = toHash(name); + detail::curve::Curve::getInstance().unregisterVariableByHash(cvh); // Update free space // Remove from properties map auto realprop = properties.find(name); @@ -662,22 +662,22 @@ void EnvironmentManager::updateDevice(const unsigned int &instance_id) { rtc_update_required = false; } if (curve_registration_required) { - auto &curve = Curve::getInstance(); + auto &curve = detail::curve::Curve::getInstance(); // Update cub for any not mapped properties for (auto &p : properties) { if (p.first.first == instance_id) { // Generate hash for the subproperty name - Curve::VariableHash cvh = toHash(p.first); + detail::curve::Curve::VariableHash cvh = toHash(p.first); const auto &prop = p.second; // Create the mapping const auto CURVE_RESULT = curve.registerVariableByHash(cvh, reinterpret_cast(prop.offset), prop.length / prop.elements, prop.elements); - if (CURVE_RESULT == Curve::UNKNOWN_VARIABLE) { + if (CURVE_RESULT == detail::curve::Curve::UNKNOWN_VARIABLE) { THROW exception::CurveException("curveRegisterVariableByHash() returned UNKNOWN_CURVE_VARIABLE, " "in EnvironmentManager::updateDevice()."); } #ifdef _DEBUG - if (CURVE_RESULT != static_cast(cvh%Curve::MAX_VARIABLES)) { + if (CURVE_RESULT != static_cast(cvh%detail::curve::Curve::MAX_VARIABLES)) { fprintf(stderr, "Curve Warning: Environment Property '%s' has a collision and may work improperly.\n", p.first.second.c_str()); } #endif @@ -687,17 +687,17 @@ void EnvironmentManager::updateDevice(const unsigned int &instance_id) { for (auto &mp : mapped_properties) { if (mp.first.first == instance_id) { // Generate hash for the subproperty name - Curve::VariableHash cvh = toHash(mp.first); + detail::curve::Curve::VariableHash cvh = toHash(mp.first); const auto &masterprop = properties.at(mp.second.masterProp); // Create the mapping const auto CURVE_RESULT = curve.registerVariableByHash(cvh, reinterpret_cast(masterprop.offset), masterprop.length / masterprop.elements, masterprop.elements); - if (CURVE_RESULT == Curve::UNKNOWN_VARIABLE) { + if (CURVE_RESULT == detail::curve::Curve::UNKNOWN_VARIABLE) { THROW exception::CurveException("curveRegisterVariableByHash() returned UNKNOWN_CURVE_VARIABLE, " "in EnvironmentManager::updateDevice()."); } #ifdef _DEBUG - if (CURVE_RESULT != static_cast(cvh%Curve::MAX_VARIABLES)) { + if (CURVE_RESULT != static_cast(cvh%detail::curve::Curve::MAX_VARIABLES)) { fprintf(stderr, "Curve Warning: Environment Property '%s' has a collision and may work improperly.\n", mp.first.second.c_str()); } #endif