.. _program_listing_file_include_flamegpu_runtime_AgentFunctionCondition.cuh: Program Listing for File AgentFunctionCondition.cuh =================================================== |exhale_lsh| :ref:`Return to documentation for file ` (``include/flamegpu/runtime/AgentFunctionCondition.cuh``) .. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS .. code-block:: cpp #ifndef INCLUDE_FLAMEGPU_RUNTIME_AGENTFUNCTIONCONDITION_CUH_ #define INCLUDE_FLAMEGPU_RUNTIME_AGENTFUNCTIONCONDITION_CUH_ #include #include #include "flamegpu/runtime/DeviceAPI.cuh" #include "flamegpu/runtime/AgentFunctionCondition_shim.cuh" #ifndef __CUDACC_RTC__ #include "flamegpu/runtime/detail/curve/DeviceCurve.cuh" #endif namespace flamegpu { // ! FLAMEGPU function return type typedef void(AgentFunctionConditionWrapper)( #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 const unsigned int popNo, detail::curandState *d_rng, unsigned int *scanFlag_conditionResult); // Can't put __global__ in a typedef template __global__ void agent_function_condition_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 const unsigned int popNo, detail::curandState *d_rng, unsigned int *scanFlag_conditionResult) { // 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 (ReadOnlyDeviceAPI::getIndex() >= popNo) return; // create a new device FLAME_GPU instance ReadOnlyDeviceAPI api = ReadOnlyDeviceAPI(d_rng); // call the user specified device function { // Negate the return value, we want false at the start of the scattered array bool conditionResult = !(AgentFunctionCondition()(&api)); // (scan flags will be processed to filter agents scanFlag_conditionResult[ReadOnlyDeviceAPI::getIndex()] = conditionResult; } } } // namespace flamegpu #endif // INCLUDE_FLAMEGPU_RUNTIME_AGENTFUNCTIONCONDITION_CUH_