.. _program_listing_file_include_flamegpu_exception_FLAMEGPUDeviceException_device.cuh:

Program Listing for File FLAMEGPUDeviceException_device.cuh
===========================================================

|exhale_lsh| :ref:`Return to documentation for file <file_include_flamegpu_exception_FLAMEGPUDeviceException_device.cuh>` (``include/flamegpu/exception/FLAMEGPUDeviceException_device.cuh``)

.. |exhale_lsh| unicode:: U+021B0 .. UPWARDS ARROW WITH TIP LEFTWARDS

.. code-block:: cpp

   #ifndef INCLUDE_FLAMEGPU_EXCEPTION_FLAMEGPUDEVICEEXCEPTION_DEVICE_CUH_
   #define INCLUDE_FLAMEGPU_EXCEPTION_FLAMEGPUDEVICEEXCEPTION_DEVICE_CUH_
   
   #include <cuda_runtime.h>
   #include <device_launch_parameters.h>  // Required for blockIdx, when not built using nvcc / nvrtc. In which case this header file is invalid anyway. It it still required to improve MSVC intellisense though?
   
   #include <cstring>
   
   #include "flamegpu/runtime/detail/SharedBlock.h"
   
   namespace flamegpu {
   namespace exception {
   
   #if !defined(FLAMEGPU_SEATBELTS) || FLAMEGPU_SEATBELTS
   #ifdef __CUDACC__
   #define DTHROW flamegpu::exception::DeviceException::create(__FILE__, __LINE__).setMessage
   #else
   // Just trying to make host compiler happy when it sees device code by mistake
   #define DTHROW(...)
   #endif
   
   struct DeviceExceptionBuffer {
       static const unsigned int MAX_ARGS = 20;
       static const unsigned int ARG_BUFF_LEN = 4096;
       static const unsigned int FORMAT_BUFF_LEN = 4096;
       static const unsigned int FILE_BUFF_LEN = 1024;
       static const unsigned int OUT_STRING_LEN = FORMAT_BUFF_LEN * 2;
       unsigned int error_count;
       char file_path[FILE_BUFF_LEN];
       unsigned int line_no;
       unsigned int block_id[3];
       unsigned int thread_id[3];
       char format_string[FORMAT_BUFF_LEN];
       unsigned int format_args_sizes[MAX_ARGS];
       char format_args[ARG_BUFF_LEN];
       unsigned int arg_count;
       unsigned int arg_offset;
   };
   #ifdef __CUDACC__
   class DeviceException {
    public:
       __device__ static DeviceException create(const char *file, const unsigned int line) {
           return {file, line};
       }
       template<typename... Args>
       __device__ void setMessage(const char *format, Args... args) {
           using detail::sm;
           // Only the thread which first reported error gets to output
           if (hasError) {
               // Only output once
               if (sm()->device_exception->format_string[0])
                   return;
               // Copy the format string
               unsigned int eos = 0;
               for (eos = 0; eos < DeviceExceptionBuffer::FORMAT_BUFF_LEN; ++eos)
                   if (format[eos] == '\0')
                       break;
               memcpy(sm()->device_exception->format_string, format, eos * sizeof(char));
               // Process args
               subformat_recurse(sm()->device_exception, args...);
           }
       }
   
    private:
       template<typename T>
       __device__ inline void subformat(DeviceExceptionBuffer *buff, T t);
       template<typename T, typename... Args>
       __device__ void subformat_recurse(DeviceExceptionBuffer *buff, const T t, Args... args) {
           // Call subformat with T
           subformat(buff, t);
           // Recurse with the rest of the list
           subformat_recurse(buff, args...);
       }
       __device__ void subformat_recurse(DeviceExceptionBuffer *buff) { }
       __device__ unsigned int strlen(const char *c) {
           unsigned int eos = 0;
           for (eos = 0; eos < DeviceExceptionBuffer::FORMAT_BUFF_LEN; ++eos)
               if (*(c + eos) == '\0' || eos >= DeviceExceptionBuffer::FORMAT_BUFF_LEN)
                   break;
           return eos + 1;  // Include the terminating character
       }
       __device__ DeviceException(const char *file, const unsigned int line)
           : hasError(!getErrorCount()) {
           using detail::sm;
           if (hasError) {
               // Copy file location
               const size_t file_len = strlen(file);
               memcpy(sm()->device_exception->file_path, file, file_len);
               // Copy line no
               sm()->device_exception->line_no = line;
               // Copy block/thread indices
               const uint3 bid3 = blockIdx;
               memcpy(sm()->device_exception->block_id, &bid3, sizeof(unsigned int) * 3);
               const uint3 tid3 = threadIdx;
               memcpy(sm()->device_exception->thread_id, &tid3, sizeof(unsigned int) * 3);
           }
       }
       __device__ inline unsigned int getErrorCount();
       const bool hasError;
   };
   template<typename T>
   __device__ inline void DeviceException::subformat(DeviceExceptionBuffer *buff, T t) {
       if (buff->arg_count < DeviceExceptionBuffer::MAX_ARGS) {
           if (buff->arg_offset + sizeof(T) <= DeviceExceptionBuffer::ARG_BUFF_LEN) {
               // Copy arg size
               buff->format_args_sizes[buff->arg_count] = sizeof(T);
               // Copy arg value
               memcpy(buff->format_args + buff->arg_offset, &t, sizeof(T));
               // Update offsets
               ++buff->arg_count;
               buff->arg_offset += sizeof(T);
           }
       }
   }
   template<>
   __device__ inline void DeviceException::subformat(DeviceExceptionBuffer *buff, const char *t) {
       if (buff->arg_count < DeviceExceptionBuffer::MAX_ARGS) {
           const unsigned int string_length = strlen(t);
           if (buff->arg_offset + string_length <= DeviceExceptionBuffer::ARG_BUFF_LEN) {
               // Copy arg size
               buff->format_args_sizes[buff->arg_count] = string_length;
               // Copy arg value
               memcpy(buff->format_args + buff->arg_offset, t, string_length);
               // Update offsets
               ++buff->arg_count;
               buff->arg_offset += string_length;
           }
       }
   }
   __device__ unsigned int DeviceException::getErrorCount() {
       using detail::sm;
       // Are we the first exception
       return atomicInc(&sm()->device_exception->error_count, UINT_MAX);
   }
   #endif
   #else
   #define DTHROW(nop)
   #endif  // FLAMEGPU_SEATBELTS=OFF
   
   }  // namespace exception
   }  // namespace flamegpu
   
   #endif  // INCLUDE_FLAMEGPU_EXCEPTION_FLAMEGPUDEVICEEXCEPTION_DEVICE_CUH_