Define CELER_UNREACHABLE on device

Before:
```
/home/s3j/.local/src/celeritas/src/orange/surfaces/detail/SurfaceAction.hh(112): warning: calling a __host__ function("__builtin_unreachable") from a __host__ __device__ function("celeritas::detail::SurfaceAction< ::celeritas_test::CalcSenseDistance> ::operator ()") is not allowed
```

Using `__trap` didn't reduce the number of instructions, but simply
omitting the warning works.

src/base/Macros.hh

@@ -92,12 +92,19 @@
  *
  * See https://clang.llvm.org/docs/LanguageExtensions.html#builtin-unreachable
  * or https://msdn.microsoft.com/en-us/library/1b3fsfxw.aspx=
+ * or
+ * https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#__builtin_unreachable
+ *
+ * (The 'unreachable' and 'assume' compiler optimizations for CUDA are only
+ * available in API version 11.1 or higher, which is encoded as major*1000 +
+ * minor*10).
  *
  * \note This macro should not generally be used; instead, the macro \c
  * CELER_ASSERT_UNREACHABLE() defined in base/Assert.hh should be used instead
  * (to provide a more detailed error message in case the point *is* reached).
  */
-#if defined(__clang__) || defined(__GNUC__)
+#if (!defined(__CUDA_ARCH__) && (defined(__clang__) || defined(__GNUC__))) \
+    || (defined(__CUDA_ARCH__) && CUDART_VERSION >= 11010)
 #    define CELER_UNREACHABLE __builtin_unreachable()
 #elif defined(_MSC_VER)
 #    define CELER_UNREACHABLE __assume(false)