diff --git a/libcudacxx/include/cuda/__complex/get_real_imag.h b/libcudacxx/include/cuda/__complex/get_real_imag.h
index 7c79efc55a4..0b2e5070966 100644
--- a/libcudacxx/include/cuda/__complex/get_real_imag.h
+++ b/libcudacxx/include/cuda/__complex/get_real_imag.h
@@ -28,24 +28,28 @@
 
 _CCCL_BEGIN_NAMESPACE_CUDA
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp>
 [[nodiscard]] _CCCL_API constexpr _Tp __get_real(const complex<_Tp>& __c) noexcept
 {
   return __c.real();
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp>
 [[nodiscard]] _CCCL_API constexpr _Tp __get_imag(const complex<_Tp>& __c) noexcept
 {
   return __c.imag();
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp>
 [[nodiscard]] _CCCL_API constexpr _Tp __get_real(const ::cuda::std::complex<_Tp>& __c) noexcept
 {
   return __c.real();
 }
 
+_CCCL_EXEC_CHECK_DISABLE
 template <class _Tp>
 [[nodiscard]] _CCCL_API constexpr _Tp __get_imag(const ::cuda::std::complex<_Tp>& __c) noexcept
 {
diff --git a/libcudacxx/include/cuda/std/__cmath/logarithms.h b/libcudacxx/include/cuda/std/__cmath/logarithms.h
index 3f65a7e1e6c..3a5bb89001d 100644
--- a/libcudacxx/include/cuda/std/__cmath/logarithms.h
+++ b/libcudacxx/include/cuda/std/__cmath/logarithms.h
@@ -201,7 +201,7 @@ template <class _Integer, enable_if_t<is_integral_v<_Integer>, int> = 0>
 #endif // _CCCL_HAS_LONG_DOUBLE()
 
 #if _LIBCUDACXX_HAS_NVFP16()
-[[nodiscard]] _CCCL_API inline __half log10(__half __x) noexcept
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half log10(__half __x) noexcept
 {
   NV_IF_ELSE_TARGET(
     NV_PROVIDES_SM_53, (return ::hlog10(__x);), (return __float2half(::cuda::std::log10f(__half2float(__x)));))
@@ -209,7 +209,7 @@ template <class _Integer, enable_if_t<is_integral_v<_Integer>, int> = 0>
 #endif // _LIBCUDACXX_HAS_NVFP16()
 
 #if _LIBCUDACXX_HAS_NVBF16()
-[[nodiscard]] _CCCL_API inline __nv_bfloat16 log10(__nv_bfloat16 __x) noexcept
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 log10(__nv_bfloat16 __x) noexcept
 {
   NV_IF_ELSE_TARGET(
     NV_IS_DEVICE, (return ::hlog10(__x);), (return __float2bfloat16(::cuda::std::log10f(__bfloat162float(__x)));))
diff --git a/libcudacxx/include/cuda/std/__cmath/modulo.h b/libcudacxx/include/cuda/std/__cmath/modulo.h
index 7ba1f20023a..d96db8fe6a2 100644
--- a/libcudacxx/include/cuda/std/__cmath/modulo.h
+++ b/libcudacxx/include/cuda/std/__cmath/modulo.h
@@ -180,7 +180,7 @@ template <class _A1, class _A2, enable_if_t<is_arithmetic_v<_A1> && is_arithmeti
 #endif // _CCCL_HAS_LONG_DOUBLE()
 
 #if _LIBCUDACXX_HAS_NVFP16()
-[[nodiscard]] _CCCL_API inline __half modf(__half __x, __half* __y) noexcept
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half modf(__half __x, __half* __y) noexcept
 {
   const __half __integral_part = ::cuda::std::trunc(__x);
   *__y                         = __integral_part;
@@ -191,7 +191,7 @@ template <class _A1, class _A2, enable_if_t<is_arithmetic_v<_A1> && is_arithmeti
 #endif // _LIBCUDACXX_HAS_NVFP16()
 
 #if _LIBCUDACXX_HAS_NVBF16()
-[[nodiscard]] _CCCL_API inline __nv_bfloat16 modf(__nv_bfloat16 __x, __nv_bfloat16* __y) noexcept
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 modf(__nv_bfloat16 __x, __nv_bfloat16* __y) noexcept
 {
   const __nv_bfloat16 __integral_part = ::cuda::std::trunc(__x);
   *__y                                = __integral_part;
diff --git a/libcudacxx/include/cuda/std/__cmath/rounding_functions.h b/libcudacxx/include/cuda/std/__cmath/rounding_functions.h
index 0bd1a5c3cb3..543878a12cf 100644
--- a/libcudacxx/include/cuda/std/__cmath/rounding_functions.h
+++ b/libcudacxx/include/cuda/std/__cmath/rounding_functions.h
@@ -802,14 +802,14 @@ _CCCL_API inline double nexttoward(_Integer __x, long double __y) noexcept
 #endif // _CCCL_HAS_LONG_DOUBLE()
 
 #if _LIBCUDACXX_HAS_NVFP16()
-[[nodiscard]] _CCCL_API inline __half rint(__half __x) noexcept
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half rint(__half __x) noexcept
 {
   NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::hrint(__x);), (return __float2half(::cuda::std::rint(__half2float(__x)));))
 }
 #endif // _LIBCUDACXX_HAS_NVFP16()
 
 #if _LIBCUDACXX_HAS_NVBF16()
-[[nodiscard]] _CCCL_API inline __nv_bfloat16 rint(__nv_bfloat16 __x) noexcept
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 rint(__nv_bfloat16 __x) noexcept
 {
   NV_IF_ELSE_TARGET(
     NV_IS_DEVICE, (return ::hrint(__x);), (return __float2bfloat16(::cuda::std::rint(__bfloat162float(__x)));))
@@ -953,14 +953,14 @@ template <class _Integer, enable_if_t<is_integral_v<_Integer>, int> = 0>
 #endif // _CCCL_HAS_LONG_DOUBLE()
 
 #if _LIBCUDACXX_HAS_NVFP16()
-[[nodiscard]] _CCCL_API inline __half trunc(__half __x) noexcept
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half trunc(__half __x) noexcept
 {
   NV_IF_ELSE_TARGET(NV_IS_DEVICE, (return ::htrunc(__x);), (return __float2half(::cuda::std::trunc(__half2float(__x)));))
 }
 #endif // _LIBCUDACXX_HAS_NVFP16()
 
 #if _LIBCUDACXX_HAS_NVBF16()
-[[nodiscard]] _CCCL_API inline __nv_bfloat16 trunc(__nv_bfloat16 __x) noexcept
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 trunc(__nv_bfloat16 __x) noexcept
 {
   NV_IF_ELSE_TARGET(
     NV_IS_DEVICE, (return ::htrunc(__x);), (return __float2bfloat16(::cuda::std::trunc(__bfloat162float(__x)));))
diff --git a/libcudacxx/include/cuda/std/__complex/complex.h b/libcudacxx/include/cuda/std/__complex/complex.h
index ee959582441..0ef78fc9241 100644
--- a/libcudacxx/include/cuda/std/__complex/complex.h
+++ b/libcudacxx/include/cuda/std/__complex/complex.h
@@ -95,13 +95,16 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _LIBCUDACXX_COMPLEX_ALIGNAS complex
   _CCCL_HIDE_FROM_ABI constexpr complex& operator=(const complex&) noexcept(is_nothrow_copy_assignable_v<_Tp>) = default;
   _CCCL_HIDE_FROM_ABI constexpr complex& operator=(complex&&) noexcept(is_nothrow_move_assignable_v<_Tp>) = default;
 
-  template <class _Up, enable_if_t<__cccl_internal::__is_non_narrowing_convertible<_Tp, _Up>::value, int> = 0>
+  template <class _Up,
+            enable_if_t<!__is_extended_floating_point_v<_Up>, int>                             = 0,
+            enable_if_t<__cccl_internal::__is_non_narrowing_convertible<_Tp, _Up>::value, int> = 0>
   _CCCL_API constexpr complex(const complex<_Up>& __c)
       : __re_(static_cast<_Tp>(__c.real()))
       , __im_(static_cast<_Tp>(__c.imag()))
   {}
 
   template <class _Up,
+            enable_if_t<!__is_extended_floating_point_v<_Up>, int>                              = 0,
             enable_if_t<!__cccl_internal::__is_non_narrowing_convertible<_Tp, _Up>::value, int> = 0,
             enable_if_t<is_constructible_v<_Tp, _Up>, int>                                      = 0>
   _CCCL_API explicit constexpr complex(const complex<_Up>& __c)
@@ -109,6 +112,9 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _LIBCUDACXX_COMPLEX_ALIGNAS complex
       , __im_(static_cast<_Tp>(__c.imag()))
   {}
 
+  template <class _Up, enable_if_t<__is_extended_floating_point_v<_Up>, int> = 0>
+  _CCCL_HOST_DEVICE_API complex(const complex<_Up>& __c);
+
   _CCCL_API constexpr complex& operator=(const value_type& __re)
   {
     __re_ = __re;
@@ -116,7 +122,7 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _LIBCUDACXX_COMPLEX_ALIGNAS complex
     return *this;
   }
 
-  template <class _Up>
+  template <class _Up, enable_if_t<!__is_extended_floating_point_v<_Up>, int> = 0>
   _CCCL_API constexpr complex& operator=(const complex<_Up>& __c)
   {
     __re_ = __c.real();
@@ -124,6 +130,9 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _LIBCUDACXX_COMPLEX_ALIGNAS complex
     return *this;
   }
 
+  template <class _Up, enable_if_t<__is_extended_floating_point_v<_Up>, int> = 0>
+  _CCCL_HOST_DEVICE_API constexpr complex& operator=(const complex<_Up>& __c);
+
 #if _CCCL_HOSTED()
   template <class _Up>
   _CCCL_API inline complex(const ::std::complex<_Up>& __other)
diff --git a/libcudacxx/include/cuda/std/__complex/nvbf16.h b/libcudacxx/include/cuda/std/__complex/nvbf16.h
index 9436c427d42..11ea755f04b 100644
--- a/libcudacxx/include/cuda/std/__complex/nvbf16.h
+++ b/libcudacxx/include/cuda/std/__complex/nvbf16.h
@@ -120,11 +120,15 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _CCCL_ALIGNAS(alignof(__nv_bfloat162)) compl
   {}
 
 #  if !_CCCL_COMPILER(GCC, <, 10) // Old GCC considers those as deleted
+  _CCCL_EXEC_CHECK_DISABLE
   _CCCL_HIDE_FROM_ABI complex(const complex&) noexcept = default;
-  _CCCL_HIDE_FROM_ABI complex(complex&&) noexcept      = default;
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_HIDE_FROM_ABI complex(complex&&) noexcept = default;
 
+  _CCCL_EXEC_CHECK_DISABLE
   _CCCL_HIDE_FROM_ABI complex& operator=(const complex&) noexcept = default;
-  _CCCL_HIDE_FROM_ABI complex& operator=(complex&&) noexcept      = default;
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_HIDE_FROM_ABI complex& operator=(complex&&) noexcept = default;
 #  endif // !_CCCL_COMPILER(GCC, <, 10)
 
   template <class _Up, enable_if_t<__cccl_internal::__is_non_narrowing_convertible<value_type, _Up>::value, int> = 0>
@@ -276,6 +280,31 @@ _CCCL_HOST_DEVICE_API inline complex<double>& complex<double>::operator=(const c
   return *this;
 }
 
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 real(const complex<__nv_bfloat16>& __c) noexcept
+{
+  return __c.real();
+}
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __nv_bfloat16 imag(const complex<__nv_bfloat16>& __c) noexcept
+{
+  return __c.imag();
+}
+
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline complex<__nv_bfloat16> conj(const complex<__nv_bfloat16>& __c)
+{
+  return complex<__nv_bfloat16>(__c.real(), ::__hneg(__c.imag()));
+}
+
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline complex<__nv_bfloat16> proj(const complex<__nv_bfloat16>& __c)
+{
+  complex<__nv_bfloat16> __r = __c;
+  if (::cuda::std::isinf(__c.real()) || ::cuda::std::isinf(__c.imag()))
+  {
+    __r = complex<__nv_bfloat16>(
+      numeric_limits<__nv_bfloat16>::infinity(), ::cuda::std::copysign(::__float2bfloat16(0.0f), __c.imag()));
+  }
+  return __r;
+}
+
 template <>
 struct __get_complex_impl<__nv_bfloat16>
 {
diff --git a/libcudacxx/include/cuda/std/__complex/nvfp16.h b/libcudacxx/include/cuda/std/__complex/nvfp16.h
index 8d4ac55cf7c..b02a0c7df91 100644
--- a/libcudacxx/include/cuda/std/__complex/nvfp16.h
+++ b/libcudacxx/include/cuda/std/__complex/nvfp16.h
@@ -120,11 +120,15 @@ class _CCCL_TYPE_VISIBILITY_DEFAULT _CCCL_ALIGNAS(alignof(__half2)) complex<__ha
   {}
 
 #  if !_CCCL_COMPILER(GCC, <, 10) // Old GCC considers those as deleted
+  _CCCL_EXEC_CHECK_DISABLE
   _CCCL_HIDE_FROM_ABI complex(const complex&) noexcept = default;
-  _CCCL_HIDE_FROM_ABI complex(complex&&) noexcept      = default;
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_HIDE_FROM_ABI complex(complex&&) noexcept = default;
 
+  _CCCL_EXEC_CHECK_DISABLE
   _CCCL_HIDE_FROM_ABI complex& operator=(const complex&) noexcept = default;
-  _CCCL_HIDE_FROM_ABI complex& operator=(complex&&) noexcept      = default;
+  _CCCL_EXEC_CHECK_DISABLE
+  _CCCL_HIDE_FROM_ABI complex& operator=(complex&&) noexcept = default;
 #  endif // !_CCCL_COMPILER(GCC, <, 10)
 
   template <class _Up, enable_if_t<__cccl_internal::__is_non_narrowing_convertible<value_type, _Up>::value, int> = 0>
@@ -276,6 +280,30 @@ _CCCL_HOST_DEVICE_API inline complex<double>& complex<double>::operator=(const c
   return *this;
 }
 
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half real(const complex<__half>& __c) noexcept
+{
+  return __c.real();
+}
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline __half imag(const complex<__half>& __c) noexcept
+{
+  return __c.imag();
+}
+
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline complex<__half> conj(const complex<__half>& __c)
+{
+  return complex<__half>(__c.real(), ::__hneg(__c.imag()));
+}
+
+[[nodiscard]] _CCCL_HOST_DEVICE_API inline complex<__half> proj(const complex<__half>& __c)
+{
+  complex<__half> __r = __c;
+  if (::cuda::std::isinf(__c.real()) || ::cuda::std::isinf(__c.imag()))
+  {
+    __r = complex<__half>(numeric_limits<__half>::infinity(), ::cuda::std::copysign(::__float2half(0.0f), __c.imag()));
+  }
+  return __r;
+}
+
 template <>
 struct __get_complex_impl<__half>
 {
diff --git a/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.assign.pass.cpp b/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.assign.pass.cpp
index b013f96b503..012681caeee 100644
--- a/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.assign.pass.cpp
+++ b/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.assign.pass.cpp
@@ -16,7 +16,7 @@
 #include "test_macros.h"
 
 template <class T, class U>
-TEST_FUNC void test_assignment()
+TEST_HOST_DEVICE_FUNC void test_assignment()
 {
   cuda::std::complex<T> from_only_real{static_cast<T>(-1.0), static_cast<T>(1.0)};
   cuda::std::complex<T> from_only_imag{static_cast<T>(-1.0), static_cast<T>(1.0)};
@@ -38,7 +38,7 @@ TEST_FUNC void test_assignment()
   assert(from_real_imag.imag() == static_cast<T>(112.0));
 }
 
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
 #if _LIBCUDACXX_HAS_NVFP16()
   test_assignment<__half, float>();
diff --git a/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.bad_macros.pass.cpp b/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.bad_macros.pass.cpp
index 6f0cb0b6d20..dcb314fa97f 100644
--- a/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.bad_macros.pass.cpp
+++ b/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.bad_macros.pass.cpp
@@ -20,7 +20,7 @@
 #include "test_macros.h"
 
 template <class T, class U>
-TEST_FUNC void test_assignment(cuda::std::complex<U> v = {})
+TEST_HOST_DEVICE_FUNC void test_assignment(cuda::std::complex<U> v = {})
 {
   cuda::std::complex<T> converting(v);
 
@@ -28,7 +28,7 @@ TEST_FUNC void test_assignment(cuda::std::complex<U> v = {})
   assigning = v;
 }
 
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
 #if _LIBCUDACXX_HAS_NVFP16()
   test_assignment<__half, float>();
diff --git a/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.cons.pass.cpp b/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.cons.pass.cpp
index 32cc98fc5c3..7aadb5bb70c 100644
--- a/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.cons.pass.cpp
+++ b/libcudacxx/test/libcudacxx/cuda/complex/half_bfloat/complex.cons.pass.cpp
@@ -16,7 +16,7 @@
 #include "test_macros.h"
 
 template <class T, class U>
-TEST_FUNC void test_construction()
+TEST_HOST_DEVICE_FUNC void test_construction()
 {
   const cuda::std::complex<U> only_real{static_cast<U>(42.0), static_cast<U>(0.0)};
   const cuda::std::complex<U> only_imag{static_cast<U>(0.0), static_cast<U>(42.0)};
@@ -34,7 +34,7 @@ TEST_FUNC void test_construction()
   assert(from_real_imag.imag() == static_cast<T>(112.0));
 }
 
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
 #if _LIBCUDACXX_HAS_NVFP16()
   test_construction<__half, float>();
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/error_functions.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/error_functions.pass.cpp
index c7d1567094b..79f41c762f6 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/error_functions.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/error_functions.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -25,7 +25,7 @@ TEST_DIAG_SUPPRESS_MSVC(4244) // conversion from 'const double' to 'int', possib
 TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
 
 template <typename T>
-TEST_FUNC void test_erf(T val)
+TEST_HOST_DEVICE_FUNC void test_erf(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::erf(T{})), ret>);
@@ -75,7 +75,7 @@ TEST_FUNC void test_erf(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_erfc(T val)
+TEST_HOST_DEVICE_FUNC void test_erfc(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::erfc(T{})), ret>);
@@ -130,13 +130,13 @@ TEST_FUNC void test_erfc(T val)
 }
 
 template <typename T>
-TEST_FUNC void test(const T val)
+TEST_HOST_DEVICE_FUNC void test(const T val)
 {
   test_erf<T>(val);
   test_erfc<T>(val);
 }
 
-TEST_FUNC void test(const float val)
+TEST_HOST_DEVICE_FUNC void test(const float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/exponential_functions.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/exponential_functions.pass.cpp
index 10d101c2445..bb2250f1d60 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/exponential_functions.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/exponential_functions.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -25,7 +25,7 @@ TEST_DIAG_SUPPRESS_MSVC(4305) // 'argument': truncation from 'T' to 'float'
 TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
 
 template <typename T>
-TEST_FUNC void test_exp(T val)
+TEST_HOST_DEVICE_FUNC void test_exp(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::exp(T{})), ret>);
@@ -55,7 +55,7 @@ TEST_FUNC void test_exp(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_exp2(T val)
+TEST_HOST_DEVICE_FUNC void test_exp2(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::exp2(T{})), ret>);
@@ -89,7 +89,7 @@ TEST_FUNC void test_exp2(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_expm1(T val)
+TEST_HOST_DEVICE_FUNC void test_expm1(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::expm1(T{})), ret>);
@@ -122,7 +122,7 @@ TEST_FUNC void test_expm1(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_frexp(T val)
+TEST_HOST_DEVICE_FUNC void test_frexp(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::frexp(T{}, nullptr)), ret>);
@@ -190,7 +190,7 @@ TEST_FUNC void test_frexp(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_ldexp(T val)
+TEST_HOST_DEVICE_FUNC void test_ldexp(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::ldexp(T{}, int{})), ret>);
@@ -236,7 +236,7 @@ TEST_FUNC void test_ldexp(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_scalbln(T val)
+TEST_HOST_DEVICE_FUNC void test_scalbln(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::scalbln(T{}, long{})), ret>);
@@ -282,7 +282,7 @@ TEST_FUNC void test_scalbln(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_scalbn(T val)
+TEST_HOST_DEVICE_FUNC void test_scalbn(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::scalbn(T{}, int{})), ret>);
@@ -328,7 +328,7 @@ TEST_FUNC void test_scalbn(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_pow(T val)
+TEST_HOST_DEVICE_FUNC void test_pow(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::pow(T{}, T{})), ret>);
@@ -388,7 +388,7 @@ TEST_FUNC void test_pow(T val)
 }
 
 template <typename T>
-TEST_FUNC void test(const T val)
+TEST_HOST_DEVICE_FUNC void test(const T val)
 {
   test_exp<T>(val);
   test_exp2<T>(val);
@@ -400,7 +400,7 @@ TEST_FUNC void test(const T val)
   test_pow<T>(val);
 }
 
-TEST_FUNC void test(const float val)
+TEST_HOST_DEVICE_FUNC void test(const float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/fdim.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/fdim.pass.cpp
index d51e92710fe..31711bf1317 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/fdim.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/fdim.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -21,7 +21,7 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(T val)
+TEST_HOST_DEVICE_FUNC void test(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::fdim(val, T())), ret>);
@@ -92,7 +92,7 @@ TEST_FUNC void test(T val)
 #endif // _CCCL_HAS_LONG_DOUBLE()
 }
 
-TEST_FUNC bool test(float val)
+TEST_HOST_DEVICE_FUNC bool test(float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/fma.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/fma.pass.cpp
index 3286392b80a..da302e1d30f 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/fma.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/fma.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -25,7 +25,7 @@ TEST_DIAG_SUPPRESS_MSVC(4305) // 'argument': truncation from 'T' to 'float'
 TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
 
 template <typename T>
-TEST_FUNC void test_fma(T val)
+TEST_HOST_DEVICE_FUNC void test_fma(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::fma(T{}, T{}, T{})), ret>);
@@ -204,12 +204,12 @@ TEST_FUNC void test_fma(T val)
 }
 
 template <typename T>
-TEST_FUNC void test(const T val)
+TEST_HOST_DEVICE_FUNC void test(const T val)
 {
   test_fma<T>(val);
 }
 
-TEST_FUNC void test(const float val)
+TEST_HOST_DEVICE_FUNC void test(const float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/gamma.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/gamma.pass.cpp
index 564d9ee9ef0..be8bb67a5cf 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/gamma.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/gamma.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -25,7 +25,7 @@ TEST_DIAG_SUPPRESS_MSVC(4305) // 'argument': truncation from 'T' to 'float'
 TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
 
 template <typename T>
-TEST_FUNC void test_lgamma(T val)
+TEST_HOST_DEVICE_FUNC void test_lgamma(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::lgamma(T{})), ret>);
@@ -119,7 +119,7 @@ TEST_FUNC void test_lgamma(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_tgamma(T val)
+TEST_HOST_DEVICE_FUNC void test_tgamma(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::tgamma(T{})), ret>);
@@ -204,13 +204,13 @@ TEST_FUNC void test_tgamma(T val)
 }
 
 template <typename T>
-TEST_FUNC void test(const T val)
+TEST_HOST_DEVICE_FUNC void test(const T val)
 {
   test_lgamma<T>(val);
   test_tgamma<T>(val);
 }
 
-TEST_FUNC void test(const float val)
+TEST_HOST_DEVICE_FUNC void test(const float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/hyperbolic_functions.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/hyperbolic_functions.pass.cpp
index 9cf2acd6dff..5b211883547 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/hyperbolic_functions.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/hyperbolic_functions.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -25,7 +25,7 @@ TEST_DIAG_SUPPRESS_MSVC(4305) // 'argument': truncation from 'T' to 'float'
 TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
 
 template <typename T>
-TEST_FUNC void test_cosh(T val)
+TEST_HOST_DEVICE_FUNC void test_cosh(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::cosh(T{})), ret>);
@@ -72,7 +72,7 @@ TEST_FUNC void test_cosh(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_sinh(T val)
+TEST_HOST_DEVICE_FUNC void test_sinh(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::sinh(T{})), ret>);
@@ -119,7 +119,7 @@ TEST_FUNC void test_sinh(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_tanh(T val)
+TEST_HOST_DEVICE_FUNC void test_tanh(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::tanh(T{})), ret>);
@@ -176,14 +176,14 @@ TEST_FUNC void test_tanh(T val)
 }
 
 template <typename T>
-TEST_FUNC void test(const T val)
+TEST_HOST_DEVICE_FUNC void test(const T val)
 {
   test_cosh<T>(val);
   test_sinh<T>(val);
   test_tanh<T>(val);
 }
 
-TEST_FUNC void test(const float val)
+TEST_HOST_DEVICE_FUNC void test(const float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/hypot.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/hypot.pass.cpp
index 99e272911c8..ab0d1340770 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/hypot.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/hypot.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -25,7 +25,7 @@ TEST_DIAG_SUPPRESS_MSVC(4305) // 'argument': truncation from 'T' to 'float'
 TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
 
 template <typename T>
-TEST_FUNC void test_hypot(T val)
+TEST_HOST_DEVICE_FUNC void test_hypot(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::hypot(T{}, T{})), ret>);
@@ -134,7 +134,7 @@ TEST_FUNC void test_hypot(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_hypot3(T val)
+TEST_HOST_DEVICE_FUNC void test_hypot3(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::hypot(T{}, T{}, T{})), ret>);
@@ -218,13 +218,13 @@ TEST_FUNC void test_hypot3(T val)
 }
 
 template <typename T>
-TEST_FUNC void test(const T val)
+TEST_HOST_DEVICE_FUNC void test(const T val)
 {
   test_hypot<T>(val);
   test_hypot3<T>(val);
 }
 
-TEST_FUNC void test(const float val)
+TEST_HOST_DEVICE_FUNC void test(const float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/inverse_hyperbolic_functions.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/inverse_hyperbolic_functions.pass.cpp
index d407562cf03..baf2284eac3 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/inverse_hyperbolic_functions.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/inverse_hyperbolic_functions.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -25,7 +25,7 @@ TEST_DIAG_SUPPRESS_MSVC(4305) // 'argument': truncation from 'T' to 'float'
 TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
 
 template <typename T>
-TEST_FUNC void test_acosh(T val)
+TEST_HOST_DEVICE_FUNC void test_acosh(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::acosh(T{})), ret>);
@@ -80,7 +80,7 @@ TEST_FUNC void test_acosh(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_asinh(T val)
+TEST_HOST_DEVICE_FUNC void test_asinh(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::asinh(T{})), ret>);
@@ -157,7 +157,7 @@ TEST_FUNC void test_asinh(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_atanh(T val)
+TEST_HOST_DEVICE_FUNC void test_atanh(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::atanh(T{})), ret>);
@@ -230,14 +230,14 @@ TEST_FUNC void test_atanh(T val)
 }
 
 template <typename T>
-TEST_FUNC void test(const T val)
+TEST_HOST_DEVICE_FUNC void test(const T val)
 {
   test_acosh<T>(val);
   test_asinh<T>(val);
   test_atanh<T>(val);
 }
 
-TEST_FUNC void test(const float val)
+TEST_HOST_DEVICE_FUNC void test(const float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/inverse_trigonometric_functions.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/inverse_trigonometric_functions.pass.cpp
index df291d97a37..8a84d12985a 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/inverse_trigonometric_functions.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/inverse_trigonometric_functions.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cmath>
 
@@ -26,7 +26,7 @@ TEST_DIAG_SUPPRESS_MSVC(4305) // 'argument': truncation from 'T' to 'float'
 TEST_DIAG_SUPPRESS_MSVC(4146) // unary minus operator applied to unsigned type, result still unsigned
 
 template <typename T>
-TEST_FUNC void test_acos(T val)
+TEST_HOST_DEVICE_FUNC void test_acos(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::acos(T{})), ret>);
@@ -68,7 +68,7 @@ TEST_FUNC void test_acos(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_asin(T val)
+TEST_HOST_DEVICE_FUNC void test_asin(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::asin(T{})), ret>);
@@ -115,7 +115,7 @@ TEST_FUNC void test_asin(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_atan(T val)
+TEST_HOST_DEVICE_FUNC void test_atan(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::atan(T{})), ret>);
@@ -158,7 +158,7 @@ TEST_FUNC void test_atan(T val)
 }
 
 template <typename T>
-TEST_FUNC void test_atan2(T val)
+TEST_HOST_DEVICE_FUNC void test_atan2(T val)
 {
   using ret = cuda::std::conditional_t<cuda::std::is_integral_v<T>, double, T>;
   static_assert(cuda::std::is_same_v<decltype(cuda::std::atan2(T{}, T{})), ret>);
@@ -341,7 +341,7 @@ TEST_FUNC void test_atan2(T val)
 }
 
 template <typename T>
-TEST_FUNC void test(const T val)
+TEST_HOST_DEVICE_FUNC void test(const T val)
 {
   test_acos<T>(val);
   test_asin<T>(val);
@@ -349,7 +349,7 @@ TEST_FUNC void test(const T val)
   test_atan2<T>(val);
 }
 
-TEST_FUNC void test(const float val)
+TEST_HOST_DEVICE_FUNC void test(const float val)
 {
   test<float>(val);
   test<double>(val);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpp
index f6f0381e58c..b27ad1a44b8 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/logarithms.pass.cpp
@@ -83,12 +83,15 @@ TEST_FUNC void test(float value)
 #if _CCCL_HAS_LONG_DOUBLE()
   test<long double>(value);
 #endif // _CCCL_HAS_LONG_DOUBLE()
-#if _LIBCUDACXX_HAS_NVFP16()
+
+#if !_CCCL_TILE_COMPILATION()
+#  if _LIBCUDACXX_HAS_NVFP16()
   test<__half>(__float2half(value));
-#endif // _LIBCUDACXX_HAS_NVFP16()
-#if _LIBCUDACXX_HAS_NVBF16()
+#  endif // _LIBCUDACXX_HAS_NVFP16()
+#  if _LIBCUDACXX_HAS_NVBF16()
   test<__nv_bfloat16>(__float2bfloat16(value));
-#endif // _LIBCUDACXX_HAS_NVBF16()
+#  endif // _LIBCUDACXX_HAS_NVBF16()
+#endif // !_CCCL_TILE_COMPILATION()
 
   test<unsigned short>(static_cast<unsigned short>(value));
   test<int>(static_cast<int>(value));
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/modulo.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/modulo.pass.cpp
index fa4dfcf1179..76639215289 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/modulo.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/modulo.pass.cpp
@@ -315,12 +315,14 @@ TEST_FUNC void test(const float val)
   test<long double>();
 #endif // _CCCL_HAS_LONG_DOUBLE()
 
-#if _LIBCUDACXX_HAS_NVFP16()
+#if !_CCCL_TILE_COMPILATION()
+#  if _LIBCUDACXX_HAS_NVFP16()
   test<__half>(val);
-#endif // _LIBCUDACXX_HAS_NVFP16()
-#if _LIBCUDACXX_HAS_NVBF16()
+#  endif // _LIBCUDACXX_HAS_NVFP16()
+#  if _LIBCUDACXX_HAS_NVBF16()
   test<__nv_bfloat16>(val);
-#endif // _LIBCUDACXX_HAS_NVBF16()
+#  endif // _LIBCUDACXX_HAS_NVBF16()
+#endif // !_CCCL_TILE_COMPILATION()
 
   test<unsigned short>(static_cast<unsigned short>(val));
   test<int>(static_cast<int>(val));
diff --git a/libcudacxx/test/libcudacxx/std/numerics/c.math/rounding.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/c.math/rounding.pass.cpp
index 3dbf9aa7793..ebc66572104 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/c.math/rounding.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/c.math/rounding.pass.cpp
@@ -454,12 +454,14 @@ TEST_FUNC void test(const float val)
   test<long double>();
 #endif // _CCCL_HAS_LONG_DOUBLE()
 
-#if _LIBCUDACXX_HAS_NVFP16()
+#if !_CCCL_TILE_COMPILATION()
+#  if _LIBCUDACXX_HAS_NVFP16()
   test<__half>(val);
-#endif // _LIBCUDACXX_HAS_NVFP16()
-#if _LIBCUDACXX_HAS_NVBF16()
+#  endif // _LIBCUDACXX_HAS_NVFP16()
+#  if _LIBCUDACXX_HAS_NVBF16()
   test<__nv_bfloat16>(val);
-#endif // _LIBCUDACXX_HAS_NVBF16()
+#  endif // _LIBCUDACXX_HAS_NVBF16()
+#endif // !_CCCL_TILE_COMPILATION()
 
   test<unsigned short>(static_cast<unsigned short>(val));
   test<int>(static_cast<int>(val));
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/arg.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/arg.pass.cpp
index 9d36af56b3b..896a4d204d8 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/arg.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/arg.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -24,21 +24,21 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(T x, typename cuda::std::enable_if<cuda::std::is_integral<T>::value>::type* = 0)
+TEST_HOST_DEVICE_FUNC void test(T x, typename cuda::std::enable_if<cuda::std::is_integral<T>::value>::type* = 0)
 {
   static_assert((cuda::std::is_same<decltype(cuda::std::arg(x)), double>::value));
   assert(cuda::std::arg(x) == arg(cuda::std::complex<double>(static_cast<double>(x), 0)));
 }
 
 template <class T>
-TEST_FUNC void test(T x, typename cuda::std::enable_if<!cuda::std::is_integral<T>::value>::type* = 0)
+TEST_HOST_DEVICE_FUNC void test(T x, typename cuda::std::enable_if<!cuda::std::is_integral<T>::value>::type* = 0)
 {
   static_assert((cuda::std::is_same<decltype(cuda::std::arg(x)), T>::value));
   assert(cuda::std::arg(x) == arg(cuda::std::complex<T>(x, 0)));
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test<T>(0);
   test<T>(1);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/norm.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/norm.pass.cpp
index 6f6ead4df77..abcea9c78e9 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/norm.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/norm.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -24,21 +24,21 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(T x, typename cuda::std::enable_if<cuda::std::is_integral<T>::value>::type* = 0)
+TEST_HOST_DEVICE_FUNC void test(T x, typename cuda::std::enable_if<cuda::std::is_integral<T>::value>::type* = 0)
 {
   static_assert((cuda::std::is_same<decltype(cuda::std::norm(x)), double>::value));
   assert(cuda::std::norm(x) == norm(cuda::std::complex<double>(static_cast<double>(x), 0)));
 }
 
 template <class T>
-TEST_FUNC void test(T x, typename cuda::std::enable_if<!cuda::std::is_integral<T>::value>::type* = 0)
+TEST_HOST_DEVICE_FUNC void test(T x, typename cuda::std::enable_if<!cuda::std::is_integral<T>::value>::type* = 0)
 {
   static_assert((cuda::std::is_same<decltype(cuda::std::norm(x)), T>::value));
   assert(cuda::std::norm(x) == norm(cuda::std::complex<T>(x, 0)));
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test<T>(0);
   test<T>(1);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/pow.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/pow.pass.cpp
index 582d0466fc0..f0cf9246513 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/pow.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/pow.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -34,7 +34,7 @@
 TEST_DIAG_SUPPRESS_MSVC(4244) // conversion from 'const double' to 'int', possible loss of data
 
 template <class T, class U>
-TEST_FUNC void test(T x, const cuda::std::complex<U>& y)
+TEST_HOST_DEVICE_FUNC void test(T x, const cuda::std::complex<U>& y)
 {
   using promote_t = typename cuda::std::common_type<T, U>::type;
   static_assert((cuda::std::is_same<decltype(cuda::std::pow(x, y)), cuda::std::complex<promote_t>>::value));
@@ -42,7 +42,7 @@ TEST_FUNC void test(T x, const cuda::std::complex<U>& y)
 }
 
 template <class T, class U>
-TEST_FUNC void test(const cuda::std::complex<T>& x, U y)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& x, U y)
 {
   using promote_t = typename cuda::std::common_type<T, U>::type;
   static_assert((cuda::std::is_same<decltype(cuda::std::pow(x, y)), cuda::std::complex<promote_t>>::value));
@@ -50,23 +50,23 @@ TEST_FUNC void test(const cuda::std::complex<T>& x, U y)
 }
 
 template <class T, class U>
-TEST_FUNC void test(const cuda::std::complex<T>& x, const cuda::std::complex<U>& y)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& x, const cuda::std::complex<U>& y)
 {
   using promote_t = typename cuda::std::common_type<T, U>::type;
   assert(cuda::std::pow(x, y) == pow(cuda::std::complex<promote_t>(x), cuda::std::complex<promote_t>(y)));
 }
 
 template <class T, class U>
-TEST_FUNC void test(typename cuda::std::enable_if<cuda::std::is_integral<T>::value>::type*  = 0,
-                    typename cuda::std::enable_if<!cuda::std::is_integral<U>::value>::type* = 0)
+TEST_HOST_DEVICE_FUNC void test(typename cuda::std::enable_if<cuda::std::is_integral<T>::value>::type*  = 0,
+                                typename cuda::std::enable_if<!cuda::std::is_integral<U>::value>::type* = 0)
 {
   test(T(3), cuda::std::complex<U>(4, 5));
   test(cuda::std::complex<U>(3, 4), T(5));
 }
 
 template <class T, class U>
-TEST_FUNC void test(typename cuda::std::enable_if<!cuda::std::is_integral<T>::value>::type* = 0,
-                    typename cuda::std::enable_if<!cuda::std::is_integral<U>::value>::type* = 0)
+TEST_HOST_DEVICE_FUNC void test(typename cuda::std::enable_if<!cuda::std::is_integral<T>::value>::type* = 0,
+                                typename cuda::std::enable_if<!cuda::std::is_integral<U>::value>::type* = 0)
 {
   test(T(3), cuda::std::complex<U>(4, 5));
   test(cuda::std::complex<T>(3, 4), U(5));
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/proj.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/proj.pass.cpp
index 84d19935a94..15b66d5f2e0 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/proj.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/cmplx.over/proj.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -28,21 +28,21 @@
 TEST_DIAG_SUPPRESS_MSVC(4244) // conversion from 'const double' to 'int', possible loss of data
 
 template <class T>
-TEST_FUNC void test(T x, typename cuda::std::enable_if<cuda::std::is_integral<T>::value>::type* = 0)
+TEST_HOST_DEVICE_FUNC void test(T x, typename cuda::std::enable_if<cuda::std::is_integral<T>::value>::type* = 0)
 {
   static_assert((cuda::std::is_same<decltype(cuda::std::proj(x)), cuda::std::complex<double>>::value));
   assert(cuda::std::proj(x) == proj(cuda::std::complex<double>(x, 0)));
 }
 
 template <class T>
-TEST_FUNC void test(T x, typename cuda::std::enable_if<cuda::std::is_floating_point<T>::value>::type* = 0)
+TEST_HOST_DEVICE_FUNC void test(T x, typename cuda::std::enable_if<cuda::std::is_floating_point<T>::value>::type* = 0)
 {
   static_assert((cuda::std::is_same<decltype(cuda::std::proj(x)), cuda::std::complex<T>>::value));
   assert(cuda::std::proj(x) == proj(cuda::std::complex<T>(x, 0)));
 }
 
 template <class T>
-TEST_FUNC void test(
+TEST_HOST_DEVICE_FUNC void test(
   T x,
   typename cuda::std::enable_if<!cuda::std::is_integral<T>::value && !cuda::std::is_floating_point<T>::value>::type* = 0)
 {
@@ -51,7 +51,7 @@ TEST_FUNC void test(
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test<T>(0);
   test<T>(1);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_complex.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_complex.pass.cpp
index 2e897fba1c1..685cedfe932 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_complex.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.member.ops/divide_equal_complex.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -22,7 +22,7 @@
 TEST_DIAG_SUPPRESS_MSVC(4244) // conversion from 'const double' to 'int', possible loss of data
 
 template <class T>
-TEST_FUNC constexpr bool test()
+TEST_HOST_DEVICE_FUNC constexpr bool test()
 {
   cuda::std::complex<T> c(-4, 7.5);
   const cuda::std::complex<T> c2(1.5, 2.5);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_complex.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_complex.pass.cpp
index a5ac53f8224..ec7a16e646f 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_complex.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/complex_divide_complex.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,7 +23,7 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC constexpr bool test()
+TEST_HOST_DEVICE_FUNC constexpr bool test()
 {
   cuda::std::complex<T> lhs(-4.0, 7.5);
   cuda::std::complex<T> rhs(1.5, 2.5);
@@ -34,7 +34,7 @@ TEST_FUNC constexpr bool test()
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_divide_complex.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_divide_complex.pass.cpp
index 49efa6cef99..3293aa6fc4e 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_divide_complex.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.ops/scalar_divide_complex.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -22,7 +22,7 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC constexpr bool test()
+TEST_HOST_DEVICE_FUNC constexpr bool test()
 {
   T lhs(-8.5);
   cuda::std::complex<T> rhs(1.5, 2.5);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acos.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acos.pass.cpp
index 6d2de16ddf0..4df6cda5e38 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acos.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acos.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,20 +23,20 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(acos(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(cuda::std::numeric_limits<T>::infinity(), 1),
        cuda::std::complex<T>(0, -cuda::std::numeric_limits<T>::infinity()));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   const T pi       = cuda::std::atan2(+0., -0.);
   auto testcases   = get_testcases<T>();
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acosh.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acosh.pass.cpp
index 0115d7830b4..1a1a920a583 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acosh.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/acosh.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,20 +23,20 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(acosh(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(cuda::std::numeric_limits<T>::infinity(), 1),
        cuda::std::complex<T>(cuda::std::numeric_limits<T>::infinity(), 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   const T pi       = cuda::std::atan2(+0., -0.);
   auto testcases   = get_testcases<T>();
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asin.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asin.pass.cpp
index 36583970679..37751b28184 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asin.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asin.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(asin(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(0, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   const T pi       = cuda::std::atan2(+0., -0.);
   auto testcases   = get_testcases<T>();
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asinh.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asinh.pass.cpp
index 5dc47f9e7bd..f2a585a8c3b 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asinh.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/asinh.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(asinh(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(0, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   const T pi       = cuda::std::atan2(+0., -0.);
   auto testcases   = get_testcases<T>();
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atan.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atan.pass.cpp
index 882c3abd209..2bde961aac0 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atan.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atan.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(atan(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(0, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atanh.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atanh.pass.cpp
index f34ec12d825..cb62f8728c3 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atanh.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/atanh.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(atanh(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(0, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   const T pi       = cuda::std::atan2(+0., -0.);
   auto testcases   = get_testcases<T>();
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cos.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cos.pass.cpp
index e189c21c07d..85fe75fa6e0 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cos.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cos.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(cos(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(1, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cosh.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cosh.pass.cpp
index 87d80aa0d23..f33ea55d70e 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cosh.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/cosh.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(cosh(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(1, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/exp.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/exp.pass.cpp
index 5c921021263..93730035864 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/exp.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/exp.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(exp(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(1, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log.pass.cpp
index c9a1c63f64a..daa355d2946 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(log(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(-cuda::std::numeric_limits<T>::infinity(), 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   const T pi       = cuda::std::atan2(+0., -0.);
   auto testcases   = get_testcases<T>();
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log10.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log10.pass.cpp
index f9c03218d49..0720793f5ab 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log10.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/log10.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(log10(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(-cuda::std::numeric_limits<T>::infinity(), 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_complex.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_complex.pass.cpp
index 0eae17fb24f..c3ce9b46197 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_complex.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_complex.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,7 +23,7 @@
 #include "test_macros.h"
 
 template <class T, class U = T>
-TEST_FUNC void test(const cuda::std::complex<T>& a, const cuda::std::complex<U>& b, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& a, const cuda::std::complex<U>& b, cuda::std::complex<T> x)
 {
   static_assert(cuda::std::is_same<decltype(pow(a, b)), cuda::std::complex<T>>::value);
   cuda::std::complex<T> c = pow(a, b);
@@ -32,13 +32,13 @@ TEST_FUNC void test(const cuda::std::complex<T>& a, const cuda::std::complex<U>&
 }
 
 template <class T, class U = T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(2, 3), cuda::std::complex<U>(2, 0), cuda::std::complex<T>(-5, 12));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_scalar.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_scalar.pass.cpp
index 4425a73a692..a365513a273 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_scalar.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_complex_scalar.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,7 +23,7 @@
 #include "test_macros.h"
 
 template <class T, class U = T>
-TEST_FUNC void test(const cuda::std::complex<T>& a, const U& b, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& a, const U& b, cuda::std::complex<T> x)
 {
   static_assert(cuda::std::is_same<decltype(pow(a, b)), cuda::std::complex<T>>::value);
   cuda::std::complex<T> c = pow(a, b);
@@ -32,13 +32,13 @@ TEST_FUNC void test(const cuda::std::complex<T>& a, const U& b, cuda::std::compl
 }
 
 template <class T, class U = T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(2, 3), U(2), cuda::std::complex<T>(-5, 12));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_scalar_complex.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_scalar_complex.pass.cpp
index 45f2720db1d..9f7874c9813 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_scalar_complex.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/pow_scalar_complex.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,7 +23,7 @@
 #include "test_macros.h"
 
 template <class T, class U = T>
-TEST_FUNC void test(const U& a, const cuda::std::complex<T>& b, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const U& a, const cuda::std::complex<T>& b, cuda::std::complex<T> x)
 {
   static_assert(cuda::std::is_same<decltype(pow(a, b)), cuda::std::complex<T>>::value);
   cuda::std::complex<T> c = pow(a, b);
@@ -32,13 +32,13 @@ TEST_FUNC void test(const U& a, const cuda::std::complex<T>& b, cuda::std::compl
 }
 
 template <class T, class U = T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(U(2), cuda::std::complex<T>(2), cuda::std::complex<T>(4));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sin.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sin.pass.cpp
index b0d51e31819..0518edf66ab 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sin.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sin.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(sin(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(0, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sinh.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sinh.pass.cpp
index 911882d8a49..f8bd6979c9e 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sinh.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sinh.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(sinh(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(0, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sqrt.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sqrt.pass.cpp
index c68db2bb7a8..58a955e1870 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sqrt.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/sqrt.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,7 +23,7 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   cuda::std::complex<T> a = sqrt(c);
   is_about(real(a), real(x));
@@ -31,13 +31,13 @@ TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(64, 0), cuda::std::complex<T>(8, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tan.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tan.pass.cpp
index 78a65afb90a..55e3c7744b5 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tan.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tan.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,20 +23,20 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(tan(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(0, 0));
   test(cuda::std::complex<T>(10000, -10000), cuda::std::complex<T>(0, -1));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tanh.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tanh.pass.cpp
index e541db8756c..b803cdcb010 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tanh.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.transcendentals/tanh.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const cuda::std::complex<T>& c, cuda::std::complex<T> x)
 {
   assert(tanh(c) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(cuda::std::complex<T>(0, 0), cuda::std::complex<T>(0, 0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/abs.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/abs.pass.cpp
index ce495fe2b7f..6e5d875a472 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/abs.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/abs.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,14 +23,14 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   cuda::std::complex<T> z(3, 4);
   assert(abs(z) == T(5));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/arg.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/arg.pass.cpp
index 5a104dd133d..b290c447336 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/arg.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/arg.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,14 +23,14 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   cuda::std::complex<T> z(1, 0);
   assert(arg(z) == T(0));
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   const T pi       = cuda::std::atan2(+0., -0.);
   auto testcases   = get_testcases<T>();
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/norm.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/norm.pass.cpp
index 024d44dfe2f..2ab5232aaa0 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/norm.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/norm.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,7 +23,7 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC constexpr bool test()
+TEST_HOST_DEVICE_FUNC constexpr bool test()
 {
   cuda::std::complex<T> z(3, 4);
   assert(norm(z) == T(25));
@@ -32,7 +32,7 @@ TEST_FUNC constexpr bool test()
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/polar.pass.cpp b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/polar.pass.cpp
index f1da8b5d26e..d21a0861acd 100644
--- a/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/polar.pass.cpp
+++ b/libcudacxx/test/libcudacxx/std/numerics/complex.number/complex.value.ops/polar.pass.cpp
@@ -7,8 +7,8 @@
 //
 //===----------------------------------------------------------------------===//
 
-// XFAIL: enable-tile
-// nvbug6077402: error: "call to non-tile function not supported!"
+// UNSUPPORTED: force-tile
+// error calling a __host__ __device__ function from a __host__ __device__ __tile__ function is not allowed
 
 // <cuda/std/complex>
 
@@ -23,19 +23,19 @@
 #include "test_macros.h"
 
 template <class T>
-TEST_FUNC void test(const T& rho, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const T& rho, cuda::std::complex<T> x)
 {
   assert(cuda::std::polar(rho) == x);
 }
 
 template <class T>
-TEST_FUNC void test(const T& rho, const T& theta, cuda::std::complex<T> x)
+TEST_HOST_DEVICE_FUNC void test(const T& rho, const T& theta, cuda::std::complex<T> x)
 {
   assert(cuda::std::polar(rho, theta) == x);
 }
 
 template <class T>
-TEST_FUNC void test()
+TEST_HOST_DEVICE_FUNC void test()
 {
   test(T(0), cuda::std::complex<T>(0, 0));
   test(T(1), cuda::std::complex<T>(1, 0));
@@ -46,7 +46,7 @@ TEST_FUNC void test()
 }
 
 template <class T>
-TEST_FUNC void test_edges()
+TEST_HOST_DEVICE_FUNC void test_edges()
 {
   auto testcases   = get_testcases<T>();
   const unsigned N = sizeof(testcases) / sizeof(testcases[0]);
diff --git a/libcudacxx/test/support/test_macros.h b/libcudacxx/test/support/test_macros.h
index 2f4861d95ab..15de4475406 100644
--- a/libcudacxx/test/support/test_macros.h
+++ b/libcudacxx/test/support/test_macros.h
@@ -30,6 +30,7 @@
 #define TEST_DEVICE_FUNC      _CCCL_DEVICE
 #define TEST_TILE_FUNC        _CCCL_TILE
 #define TEST_TILE_DEVICE_FUNC _CCCL_TILE _CCCL_DEVICE
+#define TEST_HOST_DEVICE_FUNC _CCCL_HOST_DEVICE
 
 // Use the CCCL C++ dialect detection
 #define TEST_STD_VER _CCCL_STD_VER