diff --git a/.clang-format b/.clang-format
index 91e32f0f374..7d3cdeaadd2 100644
--- a/.clang-format
+++ b/.clang-format
@@ -25,7 +25,6 @@ AlwaysBreakTemplateDeclarations: Yes
 AttributeMacros: [
                   '_CCCL_ALIGNAS_TYPE',
                   '_CCCL_ALIGNAS',
-                  '_CCCL_ALWAYS_INLINE',
                   '_CCCL_CONSTEXPR_CXX14',
                   '_CCCL_CONSTEXPR_CXX17',
                   '_CCCL_CONSTEXPR_CXX20',
diff --git a/cudax/include/cuda/experimental/__async/config.cuh b/cudax/include/cuda/experimental/__async/config.cuh
index 06cb16cca8c..2b874776c3f 100644
--- a/cudax/include/cuda/experimental/__async/config.cuh
+++ b/cudax/include/cuda/experimental/__async/config.cuh
@@ -34,7 +34,7 @@ namespace cuda::experimental::__async
 #  define _CUDAX_ARTIFICIAL
 #endif
 
-#define _CUDAX_ALWAYS_INLINE _CCCL_ALWAYS_INLINE _CUDAX_ARTIFICIAL _LIBCUDACXX_NODEBUG inline
+#define _CUDAX_ALWAYS_INLINE _CCCL_FORCEINLINE _CUDAX_ARTIFICIAL _LIBCUDACXX_NODEBUG
 
 // GCC struggles with guaranteed copy elision of immovable types.
 #if defined(_CCCL_COMPILER_GCC)
diff --git a/libcudacxx/include/cuda/std/__cccl/execution_space.h b/libcudacxx/include/cuda/std/__cccl/execution_space.h
index 6afad5001d1..f48080245ff 100644
--- a/libcudacxx/include/cuda/std/__cccl/execution_space.h
+++ b/libcudacxx/include/cuda/std/__cccl/execution_space.h
@@ -27,12 +27,10 @@
 #  define _CCCL_HOST        __host__
 #  define _CCCL_DEVICE      __device__
 #  define _CCCL_HOST_DEVICE __host__ __device__
-#  define _CCCL_FORCEINLINE __forceinline__
 #else // ^^^ _CCCL_CUDA_COMPILATION ^^^ / vvv !_CCCL_CUDA_COMPILATION vvv
 #  define _CCCL_HOST
 #  define _CCCL_DEVICE
 #  define _CCCL_HOST_DEVICE
-#  define _CCCL_FORCEINLINE inline
 #endif // !_CCCL_CUDA_COMPILATION
 
 #if !defined(_CCCL_EXEC_CHECK_DISABLE)
diff --git a/libcudacxx/include/cuda/std/__cccl/visibility.h b/libcudacxx/include/cuda/std/__cccl/visibility.h
index 64c964d81fb..56b881a2086 100644
--- a/libcudacxx/include/cuda/std/__cccl/visibility.h
+++ b/libcudacxx/include/cuda/std/__cccl/visibility.h
@@ -59,15 +59,15 @@
 #endif // !_CCCL_COMPILER_NVRTC
 
 #if defined(_CCCL_COMPILER_MSVC)
-#  define _CCCL_ALWAYS_INLINE __forceinline
+#  define _CCCL_FORCEINLINE __forceinline
 #else // ^^^ _CCCL_COMPILER_MSVC ^^^ / vvv _CCCL_COMPILER_MSVC vvv
-#  define _CCCL_ALWAYS_INLINE __attribute__((__always_inline__))
+#  define _CCCL_FORCEINLINE __inline__ __attribute__((__always_inline__))
 #endif // !_CCCL_COMPILER_MSVC
 
 #if __has_attribute(exclude_from_explicit_instantiation)
 #  define _CCCL_EXCLUDE_FROM_EXPLICIT_INSTANTIATION __attribute__((exclude_from_explicit_instantiation))
 #else // ^^^ exclude_from_explicit_instantiation ^^^ / vvv !exclude_from_explicit_instantiation vvv
-// NVCC complains mightily about being unable to inline functions if we use _CCCL_ALWAYS_INLINE here
+// NVCC complains mightily about being unable to inline functions if we use _CCCL_FORCEINLINE here
 #  define _CCCL_EXCLUDE_FROM_EXPLICIT_INSTANTIATION
 #endif // !exclude_from_explicit_instantiation
 
diff --git a/thrust/thrust/detail/function.h b/thrust/thrust/detail/function.h
index 27f3cd5d75f..d9f3dca19f9 100644
--- a/thrust/thrust/detail/function.h
+++ b/thrust/thrust/detail/function.h
@@ -39,7 +39,7 @@ struct wrapped_function
 
   _CCCL_EXEC_CHECK_DISABLE
   template <typename... Ts>
-  _CCCL_FORCEINLINE _CCCL_HOST_DEVICE Result operator()(Ts&&... args) const
+  inline _CCCL_HOST_DEVICE Result operator()(Ts&&... args) const
   {
     return static_cast<Result>(m_f(thrust::raw_reference_cast(::cuda::std::forward<Ts>(args))...));
   }