|
#pragma once |
|
|
|
#include <string> |
|
|
|
#include <thrust/detail/static_assert.h> |
|
#undef THRUST_STATIC_ASSERT |
|
#undef THRUST_STATIC_ASSERT_MSG |
|
|
|
#define THRUST_STATIC_ASSERT(B) unittest::assert_static((B), __FILE__, __LINE__); |
|
#define THRUST_STATIC_ASSERT_MSG(B, msg) unittest::assert_static((B), __FILE__, __LINE__); |
|
|
|
namespace unittest |
|
{ |
|
__host__ __device__ |
|
void assert_static(bool condition, const char * filename, int lineno); |
|
} |
|
|
|
#include <thrust/device_new.h> |
|
#include <thrust/device_delete.h> |
|
|
|
#if THRUST_DEVICE_SYSTEM == THRUST_DEVICE_SYSTEM_CUDA |
|
|
|
#define ASSERT_STATIC_ASSERT(X) \ |
|
{ \ |
|
bool triggered = false; \ |
|
typedef unittest::static_assert_exception ex_t; \ |
|
thrust::device_ptr<ex_t> device_ptr = thrust::device_new<ex_t>(); \ |
|
ex_t* raw_ptr = thrust::raw_pointer_cast(device_ptr); \ |
|
::cudaMemcpyToSymbol(unittest::detail::device_exception, &raw_ptr, sizeof(ex_t*)); \ |
|
try { X; } catch (ex_t) { triggered = true; } \ |
|
if (!triggered) { \ |
|
triggered = static_cast<ex_t>(*device_ptr).triggered; \ |
|
} \ |
|
thrust::device_free(device_ptr); \ |
|
raw_ptr = NULL; \ |
|
::cudaMemcpyToSymbol(unittest::detail::device_exception, &raw_ptr, sizeof(ex_t*)); \ |
|
if (!triggered) { unittest::UnitTestFailure f; f << "[" << __FILE__ << ":" << __LINE__ << "] did not trigger a THRUST_STATIC_ASSERT"; throw f; } \ |
|
} |
|
|
|
#else |
|
|
|
#define ASSERT_STATIC_ASSERT(X) \ |
|
{ \ |
|
bool triggered = false; \ |
|
typedef unittest::static_assert_exception ex_t; \ |
|
try { X; } catch (ex_t) { triggered = true; } \ |
|
if (!triggered) { unittest::UnitTestFailure f; f << "[" << __FILE__ << ":" << __LINE__ << "] did not trigger a THRUST_STATIC_ASSERT"; throw f; } \ |
|
} |
|
|
|
#endif |
|
|
|
namespace unittest |
|
{ |
|
class static_assert_exception |
|
{ |
|
public: |
|
__host__ __device__ |
|
static_assert_exception() : triggered(false) |
|
{ |
|
} |
|
|
|
__host__ __device__ |
|
static_assert_exception(const char * filename, int lineno) |
|
: triggered(true), filename(filename), lineno(lineno) |
|
{ |
|
} |
|
|
|
bool triggered; |
|
const char * filename; |
|
int lineno; |
|
}; |
|
|
|
namespace detail |
|
{ |
|
#ifdef __clang__ |
|
__attribute__((used)) |
|
#endif |
|
__device__ static static_assert_exception* device_exception = NULL; |
|
} |
|
|
|
__host__ __device__ |
|
void assert_static(bool condition, const char * filename, int lineno) |
|
{ |
|
if (!condition) |
|
{ |
|
static_assert_exception ex(filename, lineno); |
|
|
|
#ifdef __CUDA_ARCH__ |
|
*detail::device_exception = ex; |
|
#else |
|
throw ex; |
|
#endif |
|
} |
|
} |
|
} |
|
|
|
|