Class DeviceCurve
Defined in File DeviceCurve.cuh
Class Documentation
-
class DeviceCurve
The DeviceAPI for accessing a Curve table stored on the device.
cuRVE is a C library and this singleton class acts as a mechanism to ensure that any reference to the library is handled correctly. For example multiple objects may which to request that curve is initialised. This class will ensure that this function call is only made once the first time that a cuRVEInstance is required.
Public Functions
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ char *getVariablePtr(const char (&variableName)[M], const VariableHash namespace_hash, const unsigned int offset)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getVariable(const char (&variableName)[M], const VariableHash namespace_hash, const unsigned int agent_index, const unsigned int array_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getVariable_ldg(const char (&variableName)[M], const VariableHash namespace_hash, const unsigned int agent_index, const unsigned int array_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void setVariable(const char (&variableName)[M], const VariableHash namespace_hash, const T variable, const unsigned int agent_index, const unsigned int array_index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ T getAgentVariable(const char (&variableName)[M], unsigned int index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ T getMessageVariable(const char (&variableName)[M], unsigned int index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ T getAgentVariable_ldg(const char (&variableName)[M], unsigned int index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ T getMessageVariable_ldg(const char (&variableName)[M], unsigned int index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getAgentArrayVariable(const char (&variableName)[M], unsigned int agent_index, unsigned int array_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getMessageArrayVariable(const char (&variableName)[M], unsigned int message_index, unsigned int array_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getAgentArrayVariable_ldg(const char (&variableName)[M], unsigned int agent_index, unsigned int array_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getMessageArrayVariable_ldg(const char (&variableName)[M], unsigned int message_index, unsigned int array_index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ void setAgentVariable(const char (&variableName)[M], T variable, unsigned int index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ void setMessageVariable(const char (&variableName)[M], T variable, unsigned int index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ void setNewAgentVariable(const char (&variableName)[M], T variable, unsigned int index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void setAgentArrayVariable(const char (&variableName)[M], T variable, unsigned int agent_index, unsigned int array_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void setMessageArrayVariable(const char (&variableName)[M], T variable, unsigned int message_index, unsigned int array_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void setNewAgentArrayVariable(const char (&variableName)[M], T variable, unsigned int agent_index, unsigned int array_index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ T getEnvironmentProperty(const char (&propertyName)[M])
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getEnvironmentArrayProperty(const char (&propertyName)[M], unsigned int array_index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ T getEnvironmentDirectedGraphVertexProperty(VariableHash graphHash, const char (&propertyName)[M], unsigned int vertex_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getEnvironmentDirectedGraphVertexArrayProperty(VariableHash graphHash, const char (&propertyName)[M], unsigned int vertex_index, unsigned int array_index)
-
template<typename T, unsigned int M>
__device__ __forceinline__ T getEnvironmentDirectedGraphEdgeProperty(VariableHash graphHash, const char (&propertyName)[M], unsigned int edge_index)
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getEnvironmentDirectedGraphEdgeArrayProperty(VariableHash graphHash, const char (&propertyName)[M], unsigned int edge_index, unsigned int array_index)
-
template<typename T, unsigned int I, unsigned int J, unsigned int K, unsigned int W, unsigned int M>
__device__ __forceinline__ char *getEnvironmentMacroProperty(const char (&name)[M])
-
template<unsigned int M>
__device__ __forceinline__ unsigned int getVariableCount(const char (&variableName)[M], const VariableHash namespace_hash)
Public Static Functions
-
__device__ __forceinline__ static Variable getVariableIndex(VariableHash variable_hash)
Retrieve the index of the given hash within the cuRVE hashtable.
Note
Public for better SEATBELTS errors
- Parameters:
variable_hash – A cuRVE variable string hash from variableHash.
- Returns:
The index of the specified variable within the hash table or UNKNOWN_VARIABLE on failure.
-
__device__ __forceinline__ static inline void init(const CurveTable *__restrict__ d_curve_table)
Fill the shared memory curve table from the supplied device pointer
-
template<typename T, unsigned int M>
__device__ __forceinline__ static T getAgentVariable(const char (&variableName)[M], unsigned int index) Retrieve the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
index – The index of the variable in the named variable buffer. This corresponds to the agent/message agent/message index within the agent/message population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested variable
-
template<typename T, unsigned int M>
__device__ __forceinline__ static T getMessageVariable(const char (&variableName)[M], unsigned int index) Retrieve the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
index – The index of the variable in the named variable buffer. This corresponds to the agent/message agent/message index within the agent/message population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested variable
-
template<typename T, unsigned int M>
__device__ __forceinline__ static T getAgentVariable_ldg(const char (&variableName)[M], unsigned int index) Retrieve the specified agent/message (array) variable from the cuRVE hashtable.
Note
This uses the __ldg() intrinsic to access the variable via the read-only cache.
- Parameters:
variableName – A constant char array (C string) variable name.
index – The index of the variable in the named variable buffer. This corresponds to the agent/message agent/message index within the agent/message population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested variable
-
template<typename T, unsigned int M>
__device__ __forceinline__ static T getMessageVariable_ldg(const char (&variableName)[M], unsigned int index) Retrieve the specified agent/message (array) variable from the cuRVE hashtable.
Note
This uses the __ldg() intrinsic to access the variable via the read-only cache.
- Parameters:
variableName – A constant char array (C string) variable name.
index – The index of the variable in the named variable buffer. This corresponds to the agent/message agent/message index within the agent/message population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested variable
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static T getAgentArrayVariable(const char (&variableName)[M], unsigned int variable_index, unsigned int array_index) Retrieve the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
index – The index of the variable in the named variable buffer. This corresponds to the agent/message agent/message index within the agent/message population.
array_index – The index of the element in the named variable array.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
N – Length of the array variable specified by variableName
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested variable
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static T getMessageArrayVariable(const char (&variableName)[M], unsigned int variable_index, unsigned int array_index) Retrieve the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
index – The index of the variable in the named variable buffer. This corresponds to the agent/message agent/message index within the agent/message population.
array_index – The index of the element in the named variable array.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
N – Length of the array variable specified by variableName
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested variable
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static T getAgentArrayVariable_ldg(const char (&variableName)[M], unsigned int variable_index, unsigned int array_index) Retrieve the specified agent/message (array) variable from the cuRVE hashtable.
Note
This uses the __ldg() intrinsic to access the variable via the read-only cache.
- Parameters:
variableName – A constant char array (C string) variable name.
index – The index of the variable in the named variable buffer. This corresponds to the agent/message agent/message index within the agent/message population.
array_index – The index of the element in the named variable array.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
N – Length of the array variable specified by variableName
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested variable
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static T getMessageArrayVariable_ldg(const char (&variableName)[M], unsigned int variable_index, unsigned int array_index) Retrieve the specified agent/message (array) variable from the cuRVE hashtable.
Note
This uses the __ldg() intrinsic to access the variable via the read-only cache.
- Parameters:
variableName – A constant char array (C string) variable name.
index – The index of the variable in the named variable buffer. This corresponds to the agent/message agent/message index within the agent/message population.
array_index – The index of the element in the named variable array.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
N – Length of the array variable specified by variableName
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested variable
-
template<typename T, unsigned int M>
__device__ __forceinline__ static void setAgentVariable(const char (&variableName)[M], T variable, unsigned int index) Set the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
variable – The value to be stored.
index – The index of the variable in the named variable vector. This corresponds to the agent/message agent/message/new-agent index within the agent/message/new-agent population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
-
template<typename T, unsigned int M>
__device__ __forceinline__ static void setMessageVariable(const char (&variableName)[M], T variable, unsigned int index) Set the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
variable – The value to be stored.
index – The index of the variable in the named variable vector. This corresponds to the agent/message agent/message/new-agent index within the agent/message/new-agent population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
-
template<typename T, unsigned int M>
__device__ __forceinline__ static void setNewAgentVariable(const char (&variableName)[M], T variable, unsigned int index) Set the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
variable – The value to be stored.
index – The index of the variable in the named variable vector. This corresponds to the agent/message agent/message/new-agent index within the agent/message/new-agent population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static void setAgentArrayVariable(const char (&variableName)[M], T variable, unsigned int variable_index, unsigned int array_index) Set the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
variable – The value to be stored.
index – The index of the variable in the named variable vector. This corresponds to the agent/message agent/message/new-agent index within the agent/message/new-agent population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static void setMessageArrayVariable(const char (&variableName)[M], T variable, unsigned int variable_index, unsigned int array_index) Set the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
variable – The value to be stored.
index – The index of the variable in the named variable vector. This corresponds to the agent/message agent/message/new-agent index within the agent/message/new-agent population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static void setNewAgentArrayVariable(const char (&variableName)[M], T variable, unsigned int variable_index, unsigned int array_index) Set the specified agent/message (array) variable from the cuRVE hashtable.
- Parameters:
variableName – A constant char array (C string) variable name.
variable – The value to be stored.
index – The index of the variable in the named variable vector. This corresponds to the agent/message agent/message/new-agent index within the agent/message/new-agent population.
- Template Parameters:
T – The return type requested of the variable.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
-
template<typename T, unsigned int M>
__device__ __forceinline__ static T getEnvironmentProperty(const char (&propertyName)[M]) Retrieve the specified environment (array) property from the cuRVE hashtable.
- Parameters:
propertyName – A constant char array (C string) property name.
- Template Parameters:
T – The return type requested of the property.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified property is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested property
-
template<typename T, unsigned int N = 0, unsigned int M>
__device__ __forceinline__ static T getEnvironmentArrayProperty(const char (&propertyName)[M], unsigned int array_index) Retrieve the specified environment (array) property from the cuRVE hashtable.
- Parameters:
propertyName – A constant char array (C string) property name.
array_index – The index of the element in the named variable array.
- Template Parameters:
T – The return type requested of the property.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
N – (Optional) Length of the array variable specified by variableName, available for parity with other APIs, checked if provided (flamegpu must be built with FLAMEGPU_SEATBELTS enabled for device error checking).
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified property is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested property
-
__device__ __forceinline__ static unsigned int *getEnvironmentDirectedGraphPBM(VariableHash graphHash)
Retrieve the address of the specified environment directed graph’s pbm (CSR) for edges outgoing from each vertex
- Parameters:
graphHash – The hash of the affected graph
- Throws:
exception::DeviceError – (Only when SEATBELTS==ON) If the specified graph is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested address
-
__device__ __forceinline__ static unsigned int *getEnvironmentDirectedGraphIPBM(VariableHash graphHash)
Retrieve the address of the specified environment directed graph’s ipbm (CSC) for edges incoming to each vertex
- Parameters:
graphHash – The hash of the affected graph
- Throws:
exception::DeviceError – (Only when SEATBELTS==ON) If the specified graph is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested address
-
__device__ __forceinline__ static unsigned int *getEnvironmentDirectedGraphIPBMEdges(VariableHash graphHash)
Retrieve the address of the specified environment directed graph’s ipbm’s (CSCs) edge list
- Parameters:
graphHash – The hash of the affected graph
- Throws:
exception::DeviceError – (Only when SEATBELTS==ON) If the specified graph is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested address
-
template<typename T, unsigned int M>
__device__ __forceinline__ static T getEnvironmentDirectedGraphVertexProperty(VariableHash graphHash, const char (&propertyName)[M], unsigned int vertex_index) Retrieve the specified environment directed graph’s vertex (array) property from the cuRVE hashtable.
- Parameters:
propertyName – A constant char array (C string) property name.
- Template Parameters:
T – The return type requested of the property.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when SEATBELTS==ON) If the specified property is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested property
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static T getEnvironmentDirectedGraphVertexArrayProperty(VariableHash graphHash, const char (&propertyName)[M], unsigned int vertex_index, unsigned int array_index) Retrieve the specified environment directed graph’s vertex (array) property from the cuRVE hashtable.
- Parameters:
propertyName – A constant char array (C string) property name.
array_index – The index of the element in the named variable array.
- Template Parameters:
T – The return type requested of the property.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
N – (Optional) Length of the array variable specified by variableName, available for parity with other APIs, checked if provided (flamegpu must be built with SEATBELTS enabled for device error checking).
- Throws:
exception::DeviceError – (Only when SEATBELTS==ON) If the specified property is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested property
-
template<typename T, unsigned int M>
__device__ __forceinline__ static T getEnvironmentDirectedGraphEdgeProperty(VariableHash graphHash, const char (&propertyName)[M], unsigned int edge_index) Retrieve the specified environment directed graph’s edge (array) property from the cuRVE hashtable.
- Parameters:
propertyName – A constant char array (C string) property name.
- Template Parameters:
T – The return type requested of the property.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when SEATBELTS==ON) If the specified property is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested property
-
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ static T getEnvironmentDirectedGraphEdgeArrayProperty(VariableHash graphHash, const char (&propertyName)[M], unsigned int edge_index, unsigned int array_index) Retrieve the specified environment directed graph’s edge (array) property from the cuRVE hashtable.
- Parameters:
propertyName – A constant char array (C string) property name.
array_index – The index of the element in the named variable array.
- Template Parameters:
T – The return type requested of the property.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
N – (Optional) Length of the array variable specified by variableName, available for parity with other APIs, checked if provided (flamegpu must be built with SEATBELTS enabled for device error checking).
- Throws:
exception::DeviceError – (Only when SEATBELTS==ON) If the specified property is not found in the cuRVE hashtable, or it’s details are invalid.
- Returns:
The requested property
-
template<typename T, unsigned int I = 1, unsigned int J = 1, unsigned int K = 1, unsigned int W = 1, unsigned int M>
__device__ __forceinline__ static char *getEnvironmentMacroProperty(const char (&name)[M]) Retrieve the specified environment macro property buffer’s pointer from the cuRVE hashtable.
- Parameters:
name – A constant char array (C string) environment macro property name.
- Template Parameters:
T – The type of the requested environment macro property.
I – The length of the 1st dimension of the environment macro property, default 1.
J – The length of the 2nd dimension of the environment macro property, default 1.
K – The length of the 3rd dimension of the environment macro property, default 1.
W – The length of the 4th dimension of the environment macro property, default 1.
M – The length of the string literal passed to variableName. This parameter should always be implicit, and does not need to be provided.
- Throws:
exception::DeviceError – (Only when FLAMEGPU_SEATBELTS==ON) If the specified variable is not found in the cuRVE hashtable, or it’s details are invalid.
-
template<unsigned int M>
__device__ __forceinline__ static unsigned int getVariableCount(const char (&variableName)[M], VariableHash namespace_hash) Return the number of count of vars in the specified buffer as registered to Curve (Currently only used by EnvironmentDirectedGraph SEATBELTS checks, if used by others RTCCurve will need updating too)
-
__device__ __forceinline__ static bool isAgent(const char *agent_name)
When passed an agent name, returns a boolean to confirm whether it matches the name of the current agent
This function may be useful if an agent function is shared between multiple agents
Note
The performance of this function is unlikely to be cheap unless used as part of an RTC agent function.
-
__device__ __forceinline__ static bool isState(const char *agent_state)
When passed an agent state, returns a boolean to confirm whether it matches the name of the agent input state of the current agent function
This function may be useful if an agent function is shared between multiple agent states
Note
The performance of this function is unlikely to be cheap unless used as part of an RTC agent function (whereby it can be processed at compile time).
Public Static Attributes
-
static const int MAX_VARIABLES = Curve::MAX_VARIABLES
-
static const VariableHash EMPTY_FLAG = Curve::EMPTY_FLAG
-
static const int UNKNOWN_VARIABLE = -1
-
template<typename T, unsigned int N, unsigned int M>