diff --git a/benchmarks/utilities/simulated_memory_resource.hpp b/benchmarks/utilities/simulated_memory_resource.hpp
index 489339e69..993ec5ace 100644
--- a/benchmarks/utilities/simulated_memory_resource.hpp
+++ b/benchmarks/utilities/simulated_memory_resource.hpp
@@ -80,7 +80,7 @@ class simulated_memory_resource final : public device_memory_resource {
   void* do_allocate(std::size_t bytes, cuda_stream_view) override
   {
     // NOLINTNEXTLINE(cppcoreguidelines-pro-bounds-pointer-arithmetic)
-    RMM_EXPECTS(begin_ + bytes <= end_, rmm::bad_alloc, "Simulated memory size exceeded");
+    RMM_EXPECTS(begin_ + bytes <= end_, "Simulated memory size exceeded", rmm::bad_alloc);
     auto* ptr = static_cast<void*>(begin_);
     begin_ += bytes;  // NOLINT(cppcoreguidelines-pro-bounds-pointer-arithmetic)
     return ptr;
diff --git a/include/rmm/detail/error.hpp b/include/rmm/detail/error.hpp
index 5d8b340ab..329fa7022 100644
--- a/include/rmm/detail/error.hpp
+++ b/include/rmm/detail/error.hpp
@@ -95,27 +95,27 @@ class out_of_range : public std::out_of_range {
  * RMM_EXPECTS(p != nullptr, "Unexpected null pointer");
  *
  * // throws std::runtime_error
- * RMM_EXPECTS(p != nullptr, std::runtime_error, "Unexpected nullptr");
+ * RMM_EXPECTS(p != nullptr, "Unexpected nullptr", std::runtime_error);
  * ```
  * @param[in] _condition Expression that evaluates to true or false
+ * @param[in] _what  String literal description of why the exception was
+ *     thrown, i.e. why `_condition` was expected to be true.
  * @param[in] _expection_type The exception type to throw; must inherit
  *     `std::exception`. If not specified (i.e. if only two macro
  *     arguments are provided), defaults to `rmm::logic_error`
- * @param[in] _what  String literal description of why the exception was
- *     thrown, i.e. why `_condition` was expected to be true.
  * @throw `_exception_type` if the condition evaluates to 0 (false).
  */
 #define RMM_EXPECTS(...)                                           \
   GET_RMM_EXPECTS_MACRO(__VA_ARGS__, RMM_EXPECTS_3, RMM_EXPECTS_2) \
   (__VA_ARGS__)
 #define GET_RMM_EXPECTS_MACRO(_1, _2, _3, NAME, ...) NAME
-#define RMM_EXPECTS_3(_condition, _exception_type, _reason)                       \
+#define RMM_EXPECTS_3(_condition, _reason, _exception_type)                       \
   (!!(_condition)) ? static_cast<void>(0)                                         \
                    : throw _exception_type /*NOLINT(bugprone-macro-parentheses)*/ \
   {                                                                               \
     "RMM failure at: " __FILE__ ":" RMM_STRINGIFY(__LINE__) ": " _reason          \
   }
-#define RMM_EXPECTS_2(_condition, _reason) RMM_EXPECTS_3(_condition, rmm::logic_error, _reason)
+#define RMM_EXPECTS_2(_condition, _reason) RMM_EXPECTS_3(_condition, _reason, rmm::logic_error)
 
 /**
  * @brief Indicates that an erroneous code path has been taken.
diff --git a/include/rmm/device_uvector.hpp b/include/rmm/device_uvector.hpp
index 93abfa7a7..40f5d8c5d 100644
--- a/include/rmm/device_uvector.hpp
+++ b/include/rmm/device_uvector.hpp
@@ -208,7 +208,7 @@ class device_uvector {
                          cuda_stream_view stream)
   {
     RMM_EXPECTS(
-      element_index < size(), rmm::out_of_range, "Attempt to access out of bounds element.");
+      element_index < size(), "Attempt to access out of bounds element.", rmm::out_of_range);
 
     if constexpr (std::is_same<value_type, bool>::value) {
       RMM_CUDA_TRY(
@@ -256,7 +256,7 @@ class device_uvector {
   void set_element_to_zero_async(std::size_t element_index, cuda_stream_view stream)
   {
     RMM_EXPECTS(
-      element_index < size(), rmm::out_of_range, "Attempt to access out of bounds element.");
+      element_index < size(), "Attempt to access out of bounds element.", rmm::out_of_range);
     RMM_CUDA_TRY(
       cudaMemsetAsync(element_ptr(element_index), 0, sizeof(value_type), stream.value()));
   }
@@ -311,7 +311,7 @@ class device_uvector {
   [[nodiscard]] value_type element(std::size_t element_index, cuda_stream_view stream) const
   {
     RMM_EXPECTS(
-      element_index < size(), rmm::out_of_range, "Attempt to access out of bounds element.");
+      element_index < size(), "Attempt to access out of bounds element.", rmm::out_of_range);
     value_type value;
     RMM_CUDA_TRY(cudaMemcpyAsync(
       &value, element_ptr(element_index), sizeof(value), cudaMemcpyDefault, stream.value()));
diff --git a/include/rmm/mr/device/detail/stream_ordered_memory_resource.hpp b/include/rmm/mr/device/detail/stream_ordered_memory_resource.hpp
index 69932913d..53575e5ce 100644
--- a/include/rmm/mr/device/detail/stream_ordered_memory_resource.hpp
+++ b/include/rmm/mr/device/detail/stream_ordered_memory_resource.hpp
@@ -212,8 +212,8 @@ class stream_ordered_memory_resource : public crtp<PoolResource>, public device_
 
     size = rmm::detail::align_up(size, rmm::detail::CUDA_ALLOCATION_ALIGNMENT);
     RMM_EXPECTS(size <= this->underlying().get_maximum_allocation_size(),
-                rmm::out_of_memory,
-                "Maximum allocation size exceeded");
+                "Maximum allocation size exceeded",
+                rmm::out_of_memory);
     auto const block = this->underlying().get_block(size, stream_event);
 
     RMM_LOG_TRACE("[A][stream {:p}][{}B][{:p}]",