Program Listing for File FLAMEGPUDeviceException_device.cuh

↰ Return to documentation for file (include/flamegpu/exception/FLAMEGPUDeviceException_device.cuh)

#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_