.. _program_listing_file_include_flamegpu_runtime_AgentFunction.cuh: Program Listing for File AgentFunction.cuh ========================================== |exhale_lsh| :ref:`Return to documentation for file ` (``include/flamegpu/runtime/AgentFunction.cuh``) .. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS .. code-block:: cpp #ifndef INCLUDE_FLAMEGPU_RUNTIME_AGENTFUNCTION_CUH_ #define INCLUDE_FLAMEGPU_RUNTIME_AGENTFUNCTION_CUH_ #include #include #include "flamegpu/detail/curand.cuh" #include "flamegpu/runtime/detail/SharedBlock.h" #include "flamegpu/defines.h" #include "flamegpu/exception/FLAMEGPUDeviceException.cuh" #include "flamegpu/runtime/AgentFunction_shim.cuh" #ifndef __CUDACC_RTC__ #include "flamegpu/runtime/detail/curve/DeviceCurve.cuh" #endif namespace flamegpu { // ! FLAMEGPU function return type enum AGENT_STATUS { ALIVE = 1, DEAD = 0 }; typedef void(AgentFunctionWrapper)( #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS exception::DeviceExceptionBuffer *error_buffer, #endif #ifndef __CUDACC_RTC__ const detail::curve::CurveTable *d_curve_table, const char* d_agent_name, const char* d_state_name, const char* d_env_buffer, #endif id_t *d_agent_output_nextID, const unsigned int popNo, const void *in_messagelist_metadata, const void *out_messagelist_metadata, detail::curandState *d_rng, unsigned int *scanFlag_agentDeath, unsigned int *scanFlag_messageOutput, unsigned int *scanFlag_agentOutput); // Can't put __global__ in a typedef template __global__ void agent_function_wrapper( #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS exception::DeviceExceptionBuffer *error_buffer, #endif #ifndef __CUDACC_RTC__ const detail::curve::CurveTable* __restrict__ d_curve_table, const char* d_agent_name, const char* d_state_name, const char* d_env_buffer, #endif id_t *d_agent_output_nextID, const unsigned int popNo, const void *in_messagelist_metadata, const void *out_messagelist_metadata, detail::curandState *d_rng, unsigned int *scanFlag_agentDeath, unsigned int *scanFlag_messageOutput, unsigned int *scanFlag_agentOutput) { // We place these at the start of shared memory, so we can locate it anywhere in device code without a reference using detail::sm; if (threadIdx.x == 0) { #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS sm()->device_exception = error_buffer; #endif #ifndef __CUDACC_RTC__ sm()->agent_name = d_agent_name; sm()->state_name = d_state_name; sm()->env_buffer = d_env_buffer; #endif } #ifndef __CUDACC_RTC__ detail::curve::DeviceCurve::init(d_curve_table); #endif #if defined(__CUDACC__) // @todo - This should not be required. This template should only ever be processed by a CUDA compiler. // Sync the block after Thread 0 has written to shared. __syncthreads(); #endif // __CUDACC__ // Must be terminated here, else AgentRandom has bounds issues inside DeviceAPI constructor if (DeviceAPI::getIndex() >= popNo) return; // create a new device FLAME_GPU instance DeviceAPI api = DeviceAPI( d_agent_output_nextID, d_rng, scanFlag_agentOutput, MessageIn::In(in_messagelist_metadata), MessageOut::Out(out_messagelist_metadata, scanFlag_messageOutput)); // call the user specified device function AGENT_STATUS flag = AgentFunction()(&api); if (scanFlag_agentDeath) { // (scan flags will not be processed unless agent death has been requested in model definition) scanFlag_agentDeath[DeviceAPI::getIndex()] = flag; #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS } else if (flag == DEAD) { DTHROW("Agent death must be enabled per agent function when defining the model.\n"); #endif } } } // namespace flamegpu #endif // INCLUDE_FLAMEGPU_RUNTIME_AGENTFUNCTION_CUH_