Class Curve

Class Documentation

class Curve

A cuRVE instance.

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 Types

enum DeviceError

Enumerator for GPU device error code which may be raised by CUDA kernels

Values:

enumerator DEVICE_ERROR_NO_ERRORS
enumerator DEVICE_ERROR_UNKNOWN_VARIABLE
enumerator DEVICE_ERROR_VARIABLE_DISABLED
enumerator DEVICE_ERROR_UNKNOWN_TYPE
enumerator DEVICE_ERROR_UNKNOWN_LENGTH
enum HostError

Enumerator for cuRVE host error codes which may be raised by cuRVE API function calls

Values:

enumerator ERROR_NO_ERRORS
enumerator ERROR_UNKNOWN_VARIABLE
enumerator ERROR_TOO_MANY_VARIABLES
typedef int Variable
typedef unsigned int VariableHash
typedef unsigned int NamespaceHash

Public Functions

__host__ Variable getVariableHandle(VariableHash variable_hash)

Function for getting a handle (hash table index) to a cuRVE variable from a variable string hash Function performs hash collision avoidance using linear probing.

Parameters

variable_hash – A cuRVE variable string hash from variableHash.

Returns

Variable Handle for the cuRVE variable.

__host__ Variable registerVariableByHash(VariableHash variable_hash, void *d_ptr, size_t size, unsigned int length)

Function for registering a variable by a VariableHash Registers a variable by insertion in a hash table. Recommend using the provided registerVariable template function.

Parameters
  • variable_hash – A cuRVE variable string hash from variableHash.

  • d_ptr – a pointer to the vector which holds the hashed variable of give name

Returns

Variable Handle of registered variable or UNKNOWN_VARIABLE if an error is encountered.

template<unsigned int N, typename T>
__host__ Variable registerVariable(const char (&variableName)[N], void *d_ptr, unsigned int length)

Template function for registering a constant string Registers a constant string variable name by hashing and then inserting into a hash table.

Host side template class implementation

Parameters
  • variableName – A constant char array (C string) variable name.

  • d_ptr – a pointer to the vector which holds the variable of give name

Returns

Variable Handle of registered variable or UNKNOWN_VARIABLE if an error is encountered.

__host__ int size() const

Check how many items are in the hash table

__host__ void updateDevice()

Copy host structures to device

__host__ void unregisterVariableByHash(VariableHash variable_hash)

Function for un-registering a variable by a VariableHash Un-registers a variable by removal from a hash table. Recommend using the provided unregisterVariable template function.

TODO: Does un-registering imply that other variable with collisions will no longer be found. I.e. do you need to re-register all other variable when one is removed.

Parameters

variable_hash – A cuRVE variable string hash from variableHash.

template<unsigned int N>
__host__ void unregisterVariable(const char (&variableName)[N])

Template function for un-registering a constant string Un-registers a constant string variable name by hashing and then removing from the hash table.

Parameters

variableName – A constant char array (C string) variable name.

void __host__ printLastHostError(const char *file, const char *function, const int line)

Host API function for printing the last host error Prints the last host API error using the provided source location information. The preferred method for printing is to use the curveReportLastHostError macro which inserts source location information.

Parameters
  • file – A constant string filename.

  • function – A constant string function name.

  • line – A constant integer line number.

void __host__ printErrors(const char *file, const char *function, const int line)

Host API function for printing the last host or device error Prints the last device or host API error (or both) using the provided source location information. The preferred method for printing is to use the curveReportErrors macro which inserts source location information.

Parameters
  • file – A constant string filename.

  • function – A constant string function name.

  • line – A constant integer line number.

__host__ const char *getHostErrorString(HostError error_code)

Host API function for returning a constant string error description Returns an error description given a HostError error code.

Parameters

error_code – A HostError error code.

Returns

constant A string error description

__host__ HostError getLastHostError()

Host API function for returning the last reported error code.

Returns

A HostError error code

__host__ void clearErrors()

Host API function for clearing both the device and host error codes.

__host__ unsigned int checkHowManyMappedItems()
template<unsigned int N>
__device__ __host__ __forceinline__ Curve::VariableHash variableHash(const char (&str)[N])
template<typename T>
__device__ __forceinline__ T getVariableByHash(const VariableHash variable_hash, unsigned int index)
template<typename T>
__device__ __forceinline__ T getVariableByHash_ldg(const VariableHash variable_hash, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ T getArrayVariableByHash(const VariableHash variable_hash, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N>
__device__ __forceinline__ T getArrayVariableByHash_ldg(const VariableHash variable_hash, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N>
__device__ __forceinline__ T getAgentVariable(const char (&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ T getMessageVariable(const char (&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ T getVariable(const char (&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ T getAgentVariable_ldg(const char (&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ T getMessageVariable_ldg(const char (&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ T getVariable_ldg(const char (&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getAgentArrayVariable(const char (&variableName)[M], VariableHash namespace_hash, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getArrayVariable(const char (&variableName)[M], VariableHash namespace_hash, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getAgentArrayVariable_ldg(const char (&variableName)[M], VariableHash namespace_hash, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ T getArrayVariable_ldg(const char (&variableName)[M], VariableHash namespace_hash, unsigned int agent_index, unsigned int array_index)
template<typename T>
__device__ __forceinline__ void setVariableByHash(const VariableHash variable_hash, T variable, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ void setArrayVariableByHash(const VariableHash variable_hash, T variable, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N>
__device__ __forceinline__ void setAgentVariable(const char (&variableName)[N], VariableHash namespace_hash, T variable, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ void setMessageVariable(const char (&variableName)[N], VariableHash namespace_hash, T variable, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ void setNewAgentVariable(const char (&variableName)[N], VariableHash namespace_hash, T variable, unsigned int index)
template<typename T, unsigned int N>
__device__ __forceinline__ void setVariable(const char (&variableName)[N], VariableHash namespace_hash, T variable, unsigned int index)
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void setAgentArrayVariable(const char (&variableName)[M], VariableHash namespace_hash, T variable, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void setNewAgentArrayVariable(const char (&variableName)[M], VariableHash namespace_hash, T variable, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N, unsigned int M>
__device__ __forceinline__ void setArrayVariable(const char (&variableName)[M], VariableHash namespace_hash, T variable, unsigned int agent_index, unsigned int array_index)

Public Static Functions

__host__ static VariableHash variableRuntimeHash(const char *str)

Main cuRVE variable hashing function for strings of length determined at runtime and not compile time Should only be used for registered variables as this will be much slower than the compile time alternative.

Returns

a 32 bit cuRVE string variable hash.

__host__ static VariableHash variableRuntimeHash(unsigned int num)
template<unsigned int N> static __device__ __host__ static __forceinline__ VariableHash variableHash (const char(&str)[N])

Main cuRVE variable hashing function Calls recursive hashing functions.

Returns

a 32 bit cuRVE string variable hash.

static __device__ static __forceinline__ Variable getVariable (const VariableHash variable_hash)

Returns the index of the hashed variable within the hash table

Device side class implementation

static __device__ static __forceinline__ size_t getVariableSize (const VariableHash variable_hash)

Gets the size of the cuRVE variable type given the variable hash Gets the size of the cuRVE variable type given the variable hash.

Parameters

variable_hash – A cuRVE variable string hash from VariableHash.

Returns

A size_t which is the size of the variable or 0 otherwise

static __device__ static __forceinline__ unsigned int getVariableLength (const VariableHash variable_hash)

Gets the length of the cuRVE variable given the variable hash Gets the length of the cuRVE variable given the variable hash This will be 1 unless the variable is an array.

Parameters

variable_hash – A cuRVE variable string hash from VariableHash.

Returns

An unsigned int which is the number of elements within the curve variable (1 unless it’s an array)

static __device__ static __forceinline__ void * getVariablePtrByHash (const VariableHash variable_hash, size_t offset)

Device function for getting a pointer to a variable of given name Returns a generic pointer to a variable of given name at a specific offset in bytes from the start of the variable array.

Parameters
  • variable_hash – A cuRVE variable string hash from VariableHash.

  • offset – an offset into the variable array in bytes (offset is variable index * sizeof(variable type))

Returns

A generic pointer to the variable value. Will be nullptr if there is an error.

template<typename T> static __device__ static __forceinline__ T getVariableByHash (const VariableHash variable_hash, unsigned int index)

Device function for getting a single typed value from a VariableHash at a given index Returns a single value of specified type from a variableHash using the given index position. Note: No runtime type checking is done. TODO: Add runtime type checking for debug modes.

Parameters
  • variable_hash – A cuRVE variable string hash from VariableHash.

  • index – The index of the variable in the named variable vector.

Returns

T A value of given type at the given index for the variable with the provided hash. Will return 0 if an error is raised.

template<typename T> static __device__ static __forceinline__ T getVariableByHash_ldg (const VariableHash variable_hash, unsigned int index)
template<typename T, unsigned int N> static __device__ static __forceinline__ T getArrayVariableByHash (const VariableHash variable_hash, unsigned int agent_index, unsigned int array_index)
Template Parameters
  • T – Type of variable array

  • N – Length of variable array

template<typename T, unsigned int N> static __device__ static __forceinline__ T getArrayVariableByHash_ldg (const VariableHash variable_hash, unsigned int agent_index, unsigned int array_index)
template<typename T, unsigned int N> static __device__ static __forceinline__ T getAgentVariable (const char(&variableName)[N], VariableHash namespace_hash, unsigned int index)

These methods all forward to getVariable() This allows RTC to selectively remove a branch

template<typename T, unsigned int N> static __device__ static __forceinline__ T getMessageVariable (const char(&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N> static __device__ static __forceinline__ T getAgentVariable_ldg (const char(&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N> static __device__ static __forceinline__ T getMessageVariable_ldg (const char(&variableName)[N], VariableHash namespace_hash, unsigned int index)
template<typename T, unsigned int N, unsigned int M> static __device__ static __forceinline__ T getAgentArrayVariable (const char(&variableName)[M], VariableHash namespace_hash, unsigned int variable_index, unsigned int array_index)

These methods all forward to getArrayVariable() This allows RTC to selectively remove a branch Note: Currently messages do not support array variables

template<typename T, unsigned int N, unsigned int M> static __device__ static __forceinline__ T getAgentArrayVariable_ldg (const char(&variableName)[M], VariableHash namespace_hash, unsigned int variable_index, unsigned int array_index)
template<typename T> static __device__ static __forceinline__ void setVariableByHash (const VariableHash variable_hash, T value, unsigned int index)

Device function for setting a single typed value from a VariableHash Sets a single value from a variableHash using the given index position.

Parameters
  • variable_hash – A cuRVE variable string hash from VariableHash.

  • index – The index of the variable in the named variable vector.

  • value – The typed value to set at the given index.

template<typename T, unsigned int N> static __device__ static __forceinline__ void setArrayVariableByHash (const VariableHash variable_hash, T value, unsigned int agent_index, unsigned int array_index)
Template Parameters
  • T – Type of variable array

  • N – Length of variable array

template<typename T, unsigned int N> static __device__ static __forceinline__ void setAgentVariable (const char(&variableName)[N], VariableHash namespace_hash, T variable, unsigned int index)

These methods all forward to setArrayVariable() This allows RTC to selectively remove a branch

template<typename T, unsigned int N> static __device__ static __forceinline__ void setMessageVariable (const char(&variableName)[N], VariableHash namespace_hash, T variable, unsigned int index)
template<typename T, unsigned int N> static __device__ static __forceinline__ void setNewAgentVariable (const char(&variableName)[N], VariableHash namespace_hash, T variable, unsigned int index)
template<typename T, unsigned int N, unsigned int M> static __device__ static __forceinline__ void setAgentArrayVariable (const char(&variableName)[M], VariableHash namespace_hash, T variable, unsigned int variable_index, unsigned int array_index)

These methods all forward to setArrayVariable() This allows RTC to selectively remove a branch Note: Currently messages do not support array variables

template<typename T, unsigned int N, unsigned int M> static __device__ static __forceinline__ void setNewAgentArrayVariable (const char(&variableName)[M], VariableHash namespace_hash, T variable, unsigned int variable_index, unsigned int array_index)
static __device__ static __forceinline__ void printLastDeviceError (const char *file, const char *function, const int line)

Device function for printing the last device error Prints the last device error using the provided source location information. The preferred method for printing is to use the curveReportLastDeviceError macro which inserts source location information.

Parameters
  • file – A constant string filename.

  • function – A constant string function name.

  • line – A constant integer line number.

static __device__ __host__ static const __forceinline__ char * getDeviceErrorString (DeviceError error_code)

Device API function for returning a constant string error description Returns an error description given a DeviceError error code.

Parameters

error_code – A DeviceError error code.

Returns

constant A string error description

static __device__ static __forceinline__ DeviceError getLastDeviceError ()

Device API function for returning the last reported error code.

Returns

A DeviceError error code

static Curve &getInstance()

Gets the instance.

Returns

A new instance if this is the first request for an instance otherwise an existing instance.

Public Static Attributes

static const int UNKNOWN_VARIABLE = -1
static const int MAX_VARIABLES = 1024
static const VariableHash EMPTY_FLAG = 0
static const VariableHash DELETED_FLAG = 1
static std::mutex instance_mutex

Protected Functions

Curve()

Default constructor.

Private destructor to prevent this singleton being created more than once. Classes requiring a cuRVEInstance object should instead use the getInstance() method. This ensure that curveInit is only ever called once by the program. This will initialise the internal storage used for hash tables.