diff --git a/libcudacxx/examples/concurrent_hash_table.cu b/libcudacxx/examples/concurrent_hash_table.cu
index 717cf2f034f..51e98649881 100644
--- a/libcudacxx/examples/concurrent_hash_table.cu
+++ b/libcudacxx/examples/concurrent_hash_table.cu
@@ -24,7 +24,7 @@
 
 template <typename Key,
           typename Value,
-          typename Hash           = thrust::identity<Key>,
+          typename Hash           = cuda::std::identity,
           typename KeyEqual       = thrust::equal_to<Key>,
           typename MemoryResource = thrust::universal_memory_resource>
 struct concurrent_hash_table
diff --git a/thrust/examples/counting_iterator.cu b/thrust/examples/counting_iterator.cu
index 2979359b61d..96da36af783 100644
--- a/thrust/examples/counting_iterator.cu
+++ b/thrust/examples/counting_iterator.cu
@@ -31,7 +31,7 @@ int main()
   // compute indices of nonzero elements
   using IndexIterator = thrust::device_vector<int>::iterator;
 
-  IndexIterator indices_end = thrust::copy_if(first, last, stencil.begin(), indices.begin(), thrust::identity<int>());
+  IndexIterator indices_end = thrust::copy_if(first, last, stencil.begin(), indices.begin(), ::cuda::std::identity{});
   // indices now contains [1,2,5,7]
 
   // print result
diff --git a/thrust/examples/minimal_custom_backend.cu b/thrust/examples/minimal_custom_backend.cu
index 47826c412a0..5c8fc3fba29 100644
--- a/thrust/examples/minimal_custom_backend.cu
+++ b/thrust/examples/minimal_custom_backend.cu
@@ -44,16 +44,16 @@ int main()
   my_system sys;
 
   // To invoke our version of for_each, pass sys as the first parameter
-  thrust::for_each(sys, vec.begin(), vec.end(), thrust::identity<int>());
+  thrust::for_each(sys, vec.begin(), vec.end(), ::cuda::std::negate<>{});
 
   // Other algorithms that Thrust implements with thrust::for_each will also
   // cause our version of for_each to be invoked when we pass an instance of my_system as the first parameter.
   // Even though we did not define a special version of transform, Thrust dispatches the version it knows
   // for thrust::device_execution_policy, which my_system inherits.
-  thrust::transform(sys, vec.begin(), vec.end(), vec.begin(), thrust::identity<int>());
+  thrust::transform(sys, vec.begin(), vec.end(), vec.begin(), ::cuda::std::identity{});
 
   // Invocations without my_system are handled normally.
-  thrust::for_each(vec.begin(), vec.end(), thrust::identity<int>());
+  thrust::for_each(vec.begin(), vec.end(), ::cuda::std::negate<>{});
 
   return 0;
 }
diff --git a/thrust/testing/cuda/copy_if.cu b/thrust/testing/cuda/copy_if.cu
index 5eb60c8703d..500e80bdc50 100644
--- a/thrust/testing/cuda/copy_if.cu
+++ b/thrust/testing/cuda/copy_if.cu
@@ -266,7 +266,7 @@ void TestCopyIfStencilCudaStreams(ExecutionPolicy policy)
   cudaStreamCreate(&s);
 
   Vector::iterator end =
-    thrust::copy_if(policy.on(s), data.begin(), data.end(), stencil.begin(), result.begin(), thrust::identity<T>());
+    thrust::copy_if(policy.on(s), data.begin(), data.end(), stencil.begin(), result.begin(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(end - result.begin(), 2);
   result.resize(end - result.begin());
diff --git a/thrust/testing/cuda/is_partitioned.cu b/thrust/testing/cuda/is_partitioned.cu
index 1e02ca38c28..7384b604923 100644
--- a/thrust/testing/cuda/is_partitioned.cu
+++ b/thrust/testing/cuda/is_partitioned.cu
@@ -80,23 +80,23 @@ void TestIsPartitionedCudaStreams()
 
   // empty partition
   ASSERT_EQUAL_QUIET(true,
-                     thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.begin(), thrust::identity<int>()));
+                     thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.begin(), ::cuda::std::identity{}));
 
   // one element true partition
   ASSERT_EQUAL_QUIET(
-    true, thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.begin() + 1, thrust::identity<int>()));
+    true, thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.begin() + 1, ::cuda::std::identity{}));
 
   // just true partition
   ASSERT_EQUAL_QUIET(
-    true, thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.begin() + 2, thrust::identity<int>()));
+    true, thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.begin() + 2, ::cuda::std::identity{}));
 
   // both true & false partitions
   ASSERT_EQUAL_QUIET(true,
-                     thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<int>()));
+                     thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{}));
 
   // one element false partition
   ASSERT_EQUAL_QUIET(true,
-                     thrust::is_partitioned(thrust::cuda::par.on(s), v.begin() + 3, v.end(), thrust::identity<int>()));
+                     thrust::is_partitioned(thrust::cuda::par.on(s), v.begin() + 3, v.end(), ::cuda::std::identity{}));
 
   v[0] = 1;
   v[1] = 0;
@@ -105,7 +105,7 @@ void TestIsPartitionedCudaStreams()
 
   // not partitioned
   ASSERT_EQUAL_QUIET(false,
-                     thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<int>()));
+                     thrust::is_partitioned(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{}));
 
   cudaStreamDestroy(s);
 }
diff --git a/thrust/testing/cuda/logical.cu b/thrust/testing/cuda/logical.cu
index f690fc64a03..d40726965c0 100644
--- a/thrust/testing/cuda/logical.cu
+++ b/thrust/testing/cuda/logical.cu
@@ -18,7 +18,7 @@ void TestAllOfDevice(ExecutionPolicy exec)
   thrust::device_vector<T> v(3, 1);
   thrust::device_vector<bool> result(1);
 
-  all_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity<T>(), result.begin());
+  all_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -28,7 +28,7 @@ void TestAllOfDevice(ExecutionPolicy exec)
 
   v[1] = 0;
 
-  all_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity<T>(), result.begin());
+  all_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -36,7 +36,7 @@ void TestAllOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(false, result[0]);
 
-  all_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 0, thrust::identity<T>(), result.begin());
+  all_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -44,7 +44,7 @@ void TestAllOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(true, result[0]);
 
-  all_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 1, thrust::identity<T>(), result.begin());
+  all_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -52,7 +52,7 @@ void TestAllOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(true, result[0]);
 
-  all_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 2, thrust::identity<T>(), result.begin());
+  all_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -60,7 +60,7 @@ void TestAllOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(false, result[0]);
 
-  all_of_kernel<<<1, 1>>>(exec, v.begin() + 1, v.begin() + 2, thrust::identity<T>(), result.begin());
+  all_of_kernel<<<1, 1>>>(exec, v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -92,16 +92,16 @@ void TestAllOfCudaStreams()
   cudaStream_t s;
   cudaStreamCreate(&s);
 
-  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<T>()), true);
+  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{}), true);
 
   v[1] = 0;
 
-  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{}), false);
 
-  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 0, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, thrust::identity<T>()), false);
-  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}), false);
+  ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}), false);
 
   cudaStreamDestroy(s);
 }
@@ -122,7 +122,7 @@ void TestAnyOfDevice(ExecutionPolicy exec)
   thrust::device_vector<T> v(3, 1);
   thrust::device_vector<bool> result(1);
 
-  any_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity<T>(), result.begin());
+  any_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -132,7 +132,7 @@ void TestAnyOfDevice(ExecutionPolicy exec)
 
   v[1] = 0;
 
-  any_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity<T>(), result.begin());
+  any_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -140,7 +140,7 @@ void TestAnyOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(true, result[0]);
 
-  any_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 0, thrust::identity<T>(), result.begin());
+  any_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -148,7 +148,7 @@ void TestAnyOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(false, result[0]);
 
-  any_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 1, thrust::identity<T>(), result.begin());
+  any_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -156,7 +156,7 @@ void TestAnyOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(true, result[0]);
 
-  any_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 2, thrust::identity<T>(), result.begin());
+  any_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -164,7 +164,7 @@ void TestAnyOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(true, result[0]);
 
-  any_of_kernel<<<1, 1>>>(exec, v.begin() + 1, v.begin() + 2, thrust::identity<T>(), result.begin());
+  any_of_kernel<<<1, 1>>>(exec, v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -196,16 +196,16 @@ void TestAnyOfCudaStreams()
   cudaStream_t s;
   cudaStreamCreate(&s);
 
-  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<T>()), true);
+  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{}), true);
 
   v[1] = 0;
 
-  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<T>()), true);
+  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{}), true);
 
-  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 0, thrust::identity<T>()), false);
-  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}), false);
+  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}), false);
 
   cudaStreamDestroy(s);
 }
@@ -226,7 +226,7 @@ void TestNoneOfDevice(ExecutionPolicy exec)
   thrust::device_vector<T> v(3, 1);
   thrust::device_vector<bool> result(1);
 
-  none_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity<T>(), result.begin());
+  none_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -236,7 +236,7 @@ void TestNoneOfDevice(ExecutionPolicy exec)
 
   v[1] = 0;
 
-  none_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity<T>(), result.begin());
+  none_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -244,7 +244,7 @@ void TestNoneOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(false, result[0]);
 
-  none_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 0, thrust::identity<T>(), result.begin());
+  none_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -252,7 +252,7 @@ void TestNoneOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(true, result[0]);
 
-  none_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 1, thrust::identity<T>(), result.begin());
+  none_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -260,7 +260,7 @@ void TestNoneOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(false, result[0]);
 
-  none_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 2, thrust::identity<T>(), result.begin());
+  none_of_kernel<<<1, 1>>>(exec, v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -268,7 +268,7 @@ void TestNoneOfDevice(ExecutionPolicy exec)
 
   ASSERT_EQUAL(false, result[0]);
 
-  none_of_kernel<<<1, 1>>>(exec, v.begin() + 1, v.begin() + 2, thrust::identity<T>(), result.begin());
+  none_of_kernel<<<1, 1>>>(exec, v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}, result.begin());
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -300,16 +300,16 @@ void TestNoneOfCudaStreams()
   cudaStream_t s;
   cudaStreamCreate(&s);
 
-  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{}), false);
 
   v[1] = 0;
 
-  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{}), false);
 
-  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 0, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, thrust::identity<T>()), false);
-  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, thrust::identity<T>()), false);
-  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, thrust::identity<T>()), true);
+  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}), false);
+  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}), false);
+  ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}), true);
 
   cudaStreamDestroy(s);
 }
diff --git a/thrust/testing/cuda/partition_point.cu b/thrust/testing/cuda/partition_point.cu
index 33d8dbc76d3..0c4f0d00068 100644
--- a/thrust/testing/cuda/partition_point.cu
+++ b/thrust/testing/cuda/partition_point.cu
@@ -71,11 +71,11 @@ void TestPartitionPointCudaStreams()
   cudaStream_t s;
   cudaStreamCreate(&s);
 
-  ASSERT_EQUAL_QUIET(ref, thrust::partition_point(thrust::cuda::par.on(s), first, last, thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(ref, thrust::partition_point(thrust::cuda::par.on(s), first, last, ::cuda::std::identity{}));
 
   last = v.begin() + 3;
   ref  = last;
-  ASSERT_EQUAL_QUIET(ref, thrust::partition_point(thrust::cuda::par.on(s), first, last, thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(ref, thrust::partition_point(thrust::cuda::par.on(s), first, last, ::cuda::std::identity{}));
 
   cudaStreamDestroy(s);
 }
diff --git a/thrust/testing/cuda/remove.cu b/thrust/testing/cuda/remove.cu
index b8753330eaa..97534461c81 100644
--- a/thrust/testing/cuda/remove.cu
+++ b/thrust/testing/cuda/remove.cu
@@ -405,7 +405,7 @@ void TestRemoveIfStencilCudaStreams()
   cudaStreamCreate(&s);
 
   Vector::iterator end =
-    thrust::remove_if(thrust::cuda::par.on(s), data.begin(), data.end(), stencil.begin(), thrust::identity<T>());
+    thrust::remove_if(thrust::cuda::par.on(s), data.begin(), data.end(), stencil.begin(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(end - data.begin(), 3);
   data.erase(end, data.end());
@@ -457,7 +457,7 @@ void TestRemoveCopyIfStencilCudaStreams()
   cudaStreamCreate(&s);
 
   Vector::iterator end = thrust::remove_copy_if(
-    thrust::cuda::par.on(s), data.begin(), data.end(), stencil.begin(), result.begin(), thrust::identity<T>());
+    thrust::cuda::par.on(s), data.begin(), data.end(), stencil.begin(), result.begin(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(end - result.begin(), 3);
   result.erase(end, result.end());
diff --git a/thrust/testing/cuda/tabulate.cu b/thrust/testing/cuda/tabulate.cu
index 94d07cdbc7a..bc10ae261be 100644
--- a/thrust/testing/cuda/tabulate.cu
+++ b/thrust/testing/cuda/tabulate.cu
@@ -20,7 +20,7 @@ void TestTabulateDevice(ExecutionPolicy exec)
 
   Vector v(5);
 
-  tabulate_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity<T>());
+  tabulate_kernel<<<1, 1>>>(exec, v.begin(), v.end(), ::cuda::std::identity{});
   {
     cudaError_t const err = cudaDeviceSynchronize();
     ASSERT_EQUAL(cudaSuccess, err);
@@ -72,7 +72,7 @@ void TestTabulateCudaStreams()
   cudaStream_t s;
   cudaStreamCreate(&s);
 
-  thrust::tabulate(thrust::cuda::par.on(s), v.begin(), v.end(), thrust::identity<T>());
+  thrust::tabulate(thrust::cuda::par.on(s), v.begin(), v.end(), ::cuda::std::identity{});
   cudaStreamSynchronize(s);
 
   Vector ref{0, 1, 2, 3, 4};
diff --git a/thrust/testing/cuda/transform.cu b/thrust/testing/cuda/transform.cu
index 2e474ccfb5a..72f4fc660b3 100644
--- a/thrust/testing/cuda/transform.cu
+++ b/thrust/testing/cuda/transform.cu
@@ -80,7 +80,7 @@ void TestTransformIfUnaryNoStencilDevice(ExecutionPolicy exec)
   thrust::device_vector<typename Vector::iterator> iter_vec(1);
 
   transform_if_kernel<<<1, 1>>>(
-    exec, input.begin(), input.end(), output.begin(), thrust::negate<T>(), thrust::identity<T>(), iter_vec.begin());
+    exec, input.begin(), input.end(), output.begin(), thrust::negate<T>(), ::cuda::std::identity{}, iter_vec.begin());
   cudaError_t const err = cudaDeviceSynchronize();
   ASSERT_EQUAL(cudaSuccess, err);
 
@@ -144,7 +144,7 @@ void TestTransformIfUnaryDevice(ExecutionPolicy exec)
     stencil.begin(),
     output.begin(),
     thrust::negate<T>(),
-    thrust::identity<T>(),
+    ::cuda::std::identity{},
     iter_vec.begin());
   cudaError_t const err = cudaDeviceSynchronize();
   ASSERT_EQUAL(cudaSuccess, err);
@@ -259,7 +259,7 @@ void TestTransformIfBinaryDevice(ExecutionPolicy exec)
   Vector output{1, 2, 3};
   Vector result{5, 2, -3};
 
-  thrust::identity<T> identity;
+  ::cuda::std::identity identity;
 
   thrust::device_vector<typename Vector::iterator> iter_vec(1);
 
diff --git a/thrust/testing/cuda/transform_scan.cu b/thrust/testing/cuda/transform_scan.cu
index 935d73056a6..10bcb1c96ef 100644
--- a/thrust/testing/cuda/transform_scan.cu
+++ b/thrust/testing/cuda/transform_scan.cu
@@ -288,7 +288,8 @@ void TestTransformScanConstAccumulator()
   Vector reference(5);
   Vector output(5);
 
-  thrust::transform_inclusive_scan(input.begin(), input.end(), output.begin(), thrust::identity<T>(), thrust::plus<T>());
+  thrust::transform_inclusive_scan(
+    input.begin(), input.end(), output.begin(), ::cuda::std::identity{}, thrust::plus<T>());
   thrust::inclusive_scan(input.begin(), input.end(), reference.begin(), thrust::plus<T>());
 
   ASSERT_EQUAL(output, reference);
diff --git a/thrust/testing/dereference.cu b/thrust/testing/dereference.cu
index 26d922de331..ea40829095a 100644
--- a/thrust/testing/dereference.cu
+++ b/thrust/testing/dereference.cu
@@ -63,8 +63,8 @@ void TestDeviceDereferenceTransformIterator()
   thrust::device_vector<int> input = unittest::random_integers<int>(100);
   thrust::device_vector<int> output(input.size(), 0);
 
-  simple_copy(thrust::make_transform_iterator(input.begin(), thrust::identity<int>()),
-              thrust::make_transform_iterator(input.end(), thrust::identity<int>()),
+  simple_copy(thrust::make_transform_iterator(input.begin(), ::cuda::std::identity{}),
+              thrust::make_transform_iterator(input.end(), ::cuda::std::identity{}),
               output.begin());
 
   ASSERT_EQUAL(input, output);
@@ -76,8 +76,8 @@ void TestDeviceDereferenceTransformIteratorInputConversion()
   thrust::device_vector<int> input = unittest::random_integers<int>(100);
   thrust::device_vector<double> output(input.size(), 0);
 
-  simple_copy(thrust::make_transform_iterator(input.begin(), thrust::identity<double>()),
-              thrust::make_transform_iterator(input.end(), thrust::identity<double>()),
+  simple_copy(thrust::make_transform_iterator(input.begin(), ::cuda::std::identity{}),
+              thrust::make_transform_iterator(input.end(), ::cuda::std::identity{}),
               output.begin());
 
   ASSERT_EQUAL(input == output, true);
@@ -89,8 +89,8 @@ void TestDeviceDereferenceTransformIteratorOutputConversion()
   thrust::device_vector<int> input = unittest::random_integers<int>(100);
   thrust::device_vector<double> output(input.size(), 0);
 
-  simple_copy(thrust::make_transform_iterator(input.begin(), thrust::identity<int>()),
-              thrust::make_transform_iterator(input.end(), thrust::identity<int>()),
+  simple_copy(thrust::make_transform_iterator(input.begin(), ::cuda::std::identity{}),
+              thrust::make_transform_iterator(input.end(), ::cuda::std::identity{}),
               output.begin());
 
   ASSERT_EQUAL(input == output, true);
diff --git a/thrust/testing/find.cu b/thrust/testing/find.cu
index 9901a232fc4..074802581a7 100644
--- a/thrust/testing/find.cu
+++ b/thrust/testing/find.cu
@@ -126,7 +126,7 @@ void TestFindIfDispatchExplicit()
   thrust::device_vector<int> vec(1);
 
   my_system sys(0);
-  thrust::find_if(sys, vec.begin(), vec.end(), thrust::identity<int>());
+  thrust::find_if(sys, vec.begin(), vec.end(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(true, sys.is_valid());
 }
@@ -143,7 +143,7 @@ void TestFindIfDispatchImplicit()
 {
   thrust::device_vector<int> vec(1);
 
-  thrust::find_if(thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.end()), thrust::identity<int>());
+  thrust::find_if(thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.end()), ::cuda::std::identity{});
 
   ASSERT_EQUAL(13, vec.front());
 }
@@ -177,7 +177,7 @@ void TestFindIfNotDispatchExplicit()
   thrust::device_vector<int> vec(1);
 
   my_system sys(0);
-  thrust::find_if_not(sys, vec.begin(), vec.end(), thrust::identity<int>());
+  thrust::find_if_not(sys, vec.begin(), vec.end(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(true, sys.is_valid());
 }
@@ -194,7 +194,7 @@ void TestFindIfNotDispatchImplicit()
 {
   thrust::device_vector<int> vec(1);
 
-  thrust::find_if_not(thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.end()), thrust::identity<int>());
+  thrust::find_if_not(thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.end()), ::cuda::std::identity{});
 
   ASSERT_EQUAL(13, vec.front());
 }
diff --git a/thrust/testing/functional.cu b/thrust/testing/functional.cu
index 7757ed47bed..0d15a20f610 100644
--- a/thrust/testing/functional.cu
+++ b/thrust/testing/functional.cu
@@ -190,6 +190,7 @@ typename ::cuda::std::add_const<_Tp>::type& as_const(_Tp& __t) noexcept
 // Ad-hoc testing for other functionals
 THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestIdentityFunctional()
 {
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   int i    = 42;
   double d = 3.14;
 
@@ -219,16 +220,16 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestIdentityFunctional()
   static_assert(::cuda::std::is_same<decltype(thrust::identity<int>{}(::cuda::std::move(d))), int&&>::value, "");
   static_assert(::cuda::std::is_same<decltype(thrust::identity<int>{}(static_cast<const double&&>(d))), int&&>::value,
                 "");
+  _CCCL_SUPPRESS_DEPRECATED_POP
 }
 DECLARE_UNITTEST(TestIdentityFunctional);
 
 template <class Vector>
 THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestIdentityFunctionalVector()
 {
-  using T = typename Vector::value_type;
   Vector input{0, 1, 2, 3};
   Vector output(4);
-  thrust::transform(input.begin(), input.end(), output.begin(), thrust::identity<T>());
+  thrust::transform(input.begin(), input.end(), output.begin(), ::cuda::std::identity{});
   ASSERT_EQUAL(input, output);
 }
 DECLARE_VECTOR_UNITTEST(TestIdentityFunctionalVector);
@@ -302,13 +303,11 @@ DECLARE_VECTOR_UNITTEST(TestMinimumFunctional);
 template <class Vector>
 THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestNot1()
 {
-  using T = typename Vector::value_type;
-
   Vector input{1, 0, 1, 1, 0};
 
   Vector output(5);
 
-  thrust::transform(input.begin(), input.end(), output.begin(), thrust::not_fn(thrust::identity<T>()));
+  thrust::transform(input.begin(), input.end(), output.begin(), thrust::not_fn(::cuda::std::identity{}));
 
   Vector ref{0, 1, 0, 0, 1};
   ASSERT_EQUAL(output, ref);
diff --git a/thrust/testing/is_partitioned.cu b/thrust/testing/is_partitioned.cu
index 2de28387a64..88aa9481aa1 100644
--- a/thrust/testing/is_partitioned.cu
+++ b/thrust/testing/is_partitioned.cu
@@ -16,29 +16,27 @@ struct is_even
 template <typename Vector>
 void TestIsPartitionedSimple()
 {
-  using T = typename Vector::value_type;
-
   Vector v{1, 1, 1, 0};
 
   // empty partition
-  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.begin(), thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.begin(), ::cuda::std::identity{}));
 
   // one element true partition
-  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.begin() + 1, thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.begin() + 1, ::cuda::std::identity{}));
 
   // just true partition
-  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.begin() + 2, thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.begin() + 2, ::cuda::std::identity{}));
 
   // both true & false partitions
-  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.end(), thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.end(), ::cuda::std::identity{}));
 
   // one element false partition
-  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin() + 3, v.end(), thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin() + 3, v.end(), ::cuda::std::identity{}));
 
   v = {1, 0, 1, 1};
 
   // not partitioned
-  ASSERT_EQUAL_QUIET(false, thrust::is_partitioned(v.begin(), v.end(), thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(false, thrust::is_partitioned(v.begin(), v.end(), ::cuda::std::identity{}));
 }
 DECLARE_VECTOR_UNITTEST(TestIsPartitionedSimple);
 
diff --git a/thrust/testing/logical.cu b/thrust/testing/logical.cu
index 5203708084a..083521f7d8d 100644
--- a/thrust/testing/logical.cu
+++ b/thrust/testing/logical.cu
@@ -11,16 +11,16 @@ void TestAllOf()
 
   Vector v(3, T{1});
 
-  ASSERT_EQUAL(thrust::all_of(v.begin(), v.end(), thrust::identity<T>()), true);
+  ASSERT_EQUAL(thrust::all_of(v.begin(), v.end(), ::cuda::std::identity{}), true);
 
   v[1] = T{0};
 
-  ASSERT_EQUAL(thrust::all_of(v.begin(), v.end(), thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::all_of(v.begin(), v.end(), ::cuda::std::identity{}), false);
 
-  ASSERT_EQUAL(thrust::all_of(v.begin() + 0, v.begin() + 0, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::all_of(v.begin() + 0, v.begin() + 1, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::all_of(v.begin() + 0, v.begin() + 2, thrust::identity<T>()), false);
-  ASSERT_EQUAL(thrust::all_of(v.begin() + 1, v.begin() + 2, thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::all_of(v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::all_of(v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::all_of(v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}), false);
+  ASSERT_EQUAL(thrust::all_of(v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}), false);
 }
 DECLARE_VECTOR_UNITTEST(TestAllOf);
 
@@ -66,16 +66,16 @@ void TestAnyOf()
 
   Vector v(3, T{1});
 
-  ASSERT_EQUAL(thrust::any_of(v.begin(), v.end(), thrust::identity<T>()), true);
+  ASSERT_EQUAL(thrust::any_of(v.begin(), v.end(), ::cuda::std::identity{}), true);
 
   v[1] = 0;
 
-  ASSERT_EQUAL(thrust::any_of(v.begin(), v.end(), thrust::identity<T>()), true);
+  ASSERT_EQUAL(thrust::any_of(v.begin(), v.end(), ::cuda::std::identity{}), true);
 
-  ASSERT_EQUAL(thrust::any_of(v.begin() + 0, v.begin() + 0, thrust::identity<T>()), false);
-  ASSERT_EQUAL(thrust::any_of(v.begin() + 0, v.begin() + 1, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::any_of(v.begin() + 0, v.begin() + 2, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::any_of(v.begin() + 1, v.begin() + 2, thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::any_of(v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}), false);
+  ASSERT_EQUAL(thrust::any_of(v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::any_of(v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::any_of(v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}), false);
 }
 DECLARE_VECTOR_UNITTEST(TestAnyOf);
 
@@ -121,16 +121,16 @@ void TestNoneOf()
 
   Vector v(3, T{1});
 
-  ASSERT_EQUAL(thrust::none_of(v.begin(), v.end(), thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::none_of(v.begin(), v.end(), ::cuda::std::identity{}), false);
 
   v[1] = 0;
 
-  ASSERT_EQUAL(thrust::none_of(v.begin(), v.end(), thrust::identity<T>()), false);
+  ASSERT_EQUAL(thrust::none_of(v.begin(), v.end(), ::cuda::std::identity{}), false);
 
-  ASSERT_EQUAL(thrust::none_of(v.begin() + 0, v.begin() + 0, thrust::identity<T>()), true);
-  ASSERT_EQUAL(thrust::none_of(v.begin() + 0, v.begin() + 1, thrust::identity<T>()), false);
-  ASSERT_EQUAL(thrust::none_of(v.begin() + 0, v.begin() + 2, thrust::identity<T>()), false);
-  ASSERT_EQUAL(thrust::none_of(v.begin() + 1, v.begin() + 2, thrust::identity<T>()), true);
+  ASSERT_EQUAL(thrust::none_of(v.begin() + 0, v.begin() + 0, ::cuda::std::identity{}), true);
+  ASSERT_EQUAL(thrust::none_of(v.begin() + 0, v.begin() + 1, ::cuda::std::identity{}), false);
+  ASSERT_EQUAL(thrust::none_of(v.begin() + 0, v.begin() + 2, ::cuda::std::identity{}), false);
+  ASSERT_EQUAL(thrust::none_of(v.begin() + 1, v.begin() + 2, ::cuda::std::identity{}), true);
 }
 DECLARE_VECTOR_UNITTEST(TestNoneOf);
 
diff --git a/thrust/testing/partition.cu b/thrust/testing/partition.cu
index 62cefde5975..66d638979d4 100644
--- a/thrust/testing/partition.cu
+++ b/thrust/testing/partition.cu
@@ -139,14 +139,13 @@ DECLARE_INTEGRAL_VECTOR_UNITTEST(TestStablePartitionSimple);
 template <typename Vector>
 void TestStablePartitionStencilSimple()
 {
-  using T        = typename Vector::value_type;
   using Iterator = typename Vector::iterator;
 
   Vector data{1, 2, 1, 3, 2};
 
   Vector stencil{0, 1, 0, 0, 1};
 
-  Iterator iter = thrust::stable_partition(data.begin(), data.end(), stencil.begin(), thrust::identity<T>());
+  Iterator iter = thrust::stable_partition(data.begin(), data.end(), stencil.begin(), ::cuda::std::identity{});
 
   Vector ref{2, 2, 1, 1, 3};
 
@@ -182,17 +181,14 @@ DECLARE_INTEGRAL_VECTOR_UNITTEST(TestStablePartitionCopySimple);
 template <typename Vector>
 void TestStablePartitionCopyStencilSimple()
 {
-  using T = typename Vector::value_type;
-
   Vector data{1, 2, 1, 1, 2};
-
   Vector stencil{false, true, false, false, true};
 
   Vector true_results(2);
   Vector false_results(3);
 
   thrust::pair<typename Vector::iterator, typename Vector::iterator> ends = thrust::stable_partition_copy(
-    data.begin(), data.end(), stencil.begin(), true_results.begin(), false_results.begin(), thrust::identity<T>());
+    data.begin(), data.end(), stencil.begin(), true_results.begin(), false_results.begin(), ::cuda::std::identity{});
 
   Vector true_ref(2, 2);
 
diff --git a/thrust/testing/partition_point.cu b/thrust/testing/partition_point.cu
index 61dad342ee6..5399fe76d5f 100644
--- a/thrust/testing/partition_point.cu
+++ b/thrust/testing/partition_point.cu
@@ -16,7 +16,6 @@ struct is_even
 template <typename Vector>
 void TestPartitionPointSimple()
 {
-  using T        = typename Vector::value_type;
   using Iterator = typename Vector::iterator;
 
   Vector v{1, 1, 1, 0};
@@ -25,11 +24,11 @@ void TestPartitionPointSimple()
 
   Iterator last = v.begin() + 4;
   Iterator ref  = first + 3;
-  ASSERT_EQUAL_QUIET(ref, thrust::partition_point(first, last, thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(ref, thrust::partition_point(first, last, ::cuda::std::identity{}));
 
   last = v.begin() + 3;
   ref  = last;
-  ASSERT_EQUAL_QUIET(ref, thrust::partition_point(first, last, thrust::identity<T>()));
+  ASSERT_EQUAL_QUIET(ref, thrust::partition_point(first, last, ::cuda::std::identity{}));
 }
 DECLARE_VECTOR_UNITTEST(TestPartitionPointSimple);
 
diff --git a/thrust/testing/permutation_iterator.cu b/thrust/testing/permutation_iterator.cu
index 36bc8710700..c120bbada7d 100644
--- a/thrust/testing/permutation_iterator.cu
+++ b/thrust/testing/permutation_iterator.cu
@@ -242,7 +242,7 @@ void TestPermutationIteratorWithCountingIterator()
     thrust::transform(thrust::make_permutation_iterator(input, index),
                       thrust::make_permutation_iterator(input, index + 4),
                       output.begin(),
-                      thrust::identity<T>());
+                      ::cuda::std::identity{});
 
     Vector ref{0, 1, 2, 3};
     ASSERT_EQUAL(output, ref);
diff --git a/thrust/testing/remove.cu b/thrust/testing/remove.cu
index a94f23a2ef6..53d7b60023a 100644
--- a/thrust/testing/remove.cu
+++ b/thrust/testing/remove.cu
@@ -189,13 +189,10 @@ DECLARE_UNITTEST(TestRemoveIfDispatchImplicit);
 template <typename Vector>
 void TestRemoveIfStencilSimple()
 {
-  using T = typename Vector::value_type;
-
   Vector data{1, 2, 1, 3, 2};
-
   Vector stencil{0, 1, 0, 0, 1};
 
-  typename Vector::iterator end = thrust::remove_if(data.begin(), data.end(), stencil.begin(), thrust::identity<T>());
+  typename Vector::iterator end = thrust::remove_if(data.begin(), data.end(), stencil.begin(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(end - data.begin(), 3);
   data.resize(end - data.begin());
@@ -299,16 +296,13 @@ DECLARE_UNITTEST(TestRemoveCopyIfDispatchImplicit);
 template <typename Vector>
 void TestRemoveCopyIfStencilSimple()
 {
-  using T = typename Vector::value_type;
-
   Vector data{1, 2, 1, 3, 2};
-
   Vector stencil{0, 1, 0, 0, 1};
 
   Vector result(5);
 
   typename Vector::iterator end =
-    thrust::remove_copy_if(data.begin(), data.end(), stencil.begin(), result.begin(), thrust::identity<T>());
+    thrust::remove_copy_if(data.begin(), data.end(), stencil.begin(), result.begin(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(end - result.begin(), 3);
   result.resize(end - result.begin());
diff --git a/thrust/testing/tabulate.cu b/thrust/testing/tabulate.cu
index c18007fade5..e72349ffe80 100644
--- a/thrust/testing/tabulate.cu
+++ b/thrust/testing/tabulate.cu
@@ -16,7 +16,7 @@ void TestTabulateDispatchExplicit()
   thrust::device_vector<int> vec(1);
 
   my_system sys(0);
-  thrust::tabulate(sys, vec.begin(), vec.end(), thrust::identity<int>());
+  thrust::tabulate(sys, vec.begin(), vec.end(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(true, sys.is_valid());
 }
@@ -32,7 +32,7 @@ void TestTabulateDispatchImplicit()
 {
   thrust::device_vector<int> vec(1);
 
-  thrust::tabulate(thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.end()), thrust::identity<int>());
+  thrust::tabulate(thrust::retag<my_tag>(vec.begin()), thrust::retag<my_tag>(vec.end()), ::cuda::std::identity{});
 
   ASSERT_EQUAL(13, vec.front());
 }
@@ -42,11 +42,10 @@ template <class Vector>
 void TestTabulateSimple()
 {
   using namespace thrust::placeholders;
-  using T = typename Vector::value_type;
 
   Vector v(5);
 
-  thrust::tabulate(v.begin(), v.end(), thrust::identity<T>());
+  thrust::tabulate(v.begin(), v.end(), ::cuda::std::identity{});
 
   Vector ref{0, 1, 2, 3, 4};
   ASSERT_EQUAL(v, ref);
@@ -88,7 +87,7 @@ void TestTabulateToDiscardIterator(size_t n)
 {
   thrust::tabulate(thrust::discard_iterator<thrust::device_system_tag>(),
                    thrust::discard_iterator<thrust::device_system_tag>(n),
-                   thrust::identity<int>());
+                   ::cuda::std::identity{});
 
   // nothing to check -- just make sure it compiles
 }
diff --git a/thrust/testing/transform.cu b/thrust/testing/transform.cu
index 24f035d0f9d..bc35ce4aa24 100644
--- a/thrust/testing/transform.cu
+++ b/thrust/testing/transform.cu
@@ -81,7 +81,7 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestTransformIfUnaryNoStencilSimple()
   Vector output{-1, -2, -3};
   Vector result{-1, 2, -3};
 
-  iter = thrust::transform_if(input.begin(), input.end(), output.begin(), thrust::negate<T>(), thrust::identity<T>());
+  iter = thrust::transform_if(input.begin(), input.end(), output.begin(), thrust::negate<T>(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(std::size_t(iter - output.begin()), input.size());
   ASSERT_EQUAL(output, result);
@@ -142,7 +142,7 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestTransformIfUnarySimple()
   Vector result{-1, 2, -3};
 
   iter = thrust::transform_if(
-    input.begin(), input.end(), stencil.begin(), output.begin(), thrust::negate<T>(), thrust::identity<T>());
+    input.begin(), input.end(), stencil.begin(), output.begin(), thrust::negate<T>(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(std::size_t(iter - output.begin()), input.size());
   ASSERT_EQUAL(output, result);
@@ -270,7 +270,7 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestTransformIfBinarySimple()
   Vector output{1, 2, 3};
   Vector result{5, 2, -3};
 
-  thrust::identity<T> identity;
+  ::cuda::std::identity identity;
 
   iter = thrust::transform_if(
     input1.begin(),
@@ -675,8 +675,8 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestTransformUnaryCountingIterator()
   thrust::host_vector<T> h_result(n);
   thrust::device_vector<T> d_result(n);
 
-  thrust::transform(h_first, h_first + n, h_result.begin(), thrust::identity<T>());
-  thrust::transform(d_first, d_first + n, d_result.begin(), thrust::identity<T>());
+  thrust::transform(h_first, h_first + n, h_result.begin(), ::cuda::std::identity{});
+  thrust::transform(d_first, d_first + n, d_result.begin(), ::cuda::std::identity{});
 
   ASSERT_EQUAL(h_result, d_result);
 }
diff --git a/thrust/testing/transform_input_output_iterator.cu b/thrust/testing/transform_input_output_iterator.cu
index acc5517e21f..01744301454 100644
--- a/thrust/testing/transform_input_output_iterator.cu
+++ b/thrust/testing/transform_input_output_iterator.cu
@@ -99,8 +99,8 @@ struct TestTransformInputOutputIteratorScan
 
     // run on host (uses forward iterator negate)
     thrust::inclusive_scan(
-      thrust::make_transform_input_output_iterator(h_data.begin(), thrust::negate<T>(), thrust::identity<T>()),
-      thrust::make_transform_input_output_iterator(h_data.end(), thrust::negate<T>(), thrust::identity<T>()),
+      thrust::make_transform_input_output_iterator(h_data.begin(), thrust::negate<T>(), ::cuda::std::identity{}),
+      thrust::make_transform_input_output_iterator(h_data.end(), thrust::negate<T>(), ::cuda::std::identity{}),
       h_result.begin());
     // run on device (uses reverse iterator negate)
     thrust::inclusive_scan(
diff --git a/thrust/testing/transform_iterator.cu b/thrust/testing/transform_iterator.cu
index 1c8cbb5b44f..c8e0044f99e 100644
--- a/thrust/testing/transform_iterator.cu
+++ b/thrust/testing/transform_iterator.cu
@@ -1,3 +1,10 @@
+#include <cuda/__cccl_config>
+
+#if _CCCL_COMPILER(NVHPC)
+// suppress warnings on thrust::identity
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+#endif // _CCCL_COMPILER(NVHPC)
+
 #include <thrust/copy.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/counting_iterator.h>
@@ -10,6 +17,10 @@
 
 #include <unittest/unittest.h>
 
+#if _CCCL_COMPILER(NVHPC)
+_CCCL_SUPPRESS_DEPRECATED_POP
+#endif // _CCCL_COMPILER(NVHPC)
+
 template <class Vector>
 void TestTransformIterator()
 {
@@ -134,6 +145,7 @@ struct forward
 
 void TestTransformIteratorReferenceAndValueType()
 {
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   using ::cuda::std::is_same;
   using ::cuda::std::negate;
   {
@@ -163,8 +175,9 @@ void TestTransformIteratorReferenceAndValueType()
     static_assert(is_same<decltype(it_tr_tid)::value_type, bool>::value, "");
     (void) it_tr_tid;
 
-    auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::__identity{});
-    static_assert(is_same<decltype(it_tr_cid)::reference, bool&&>::value, ""); // inferred, like forward
+    auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::identity{});
+    static_assert(is_same<decltype(it_tr_cid)::reference, bool>::value, ""); // special handling by
+                                                                             // transform_iterator_reference
     static_assert(is_same<decltype(it_tr_cid)::value_type, bool>::value, "");
     (void) it_tr_cid;
   }
@@ -196,8 +209,9 @@ void TestTransformIteratorReferenceAndValueType()
     static_assert(is_same<decltype(it_tr_tid)::value_type, bool>::value, "");
     (void) it_tr_tid;
 
-    auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::__identity{});
-    static_assert(is_same<decltype(it_tr_cid)::reference, bool&&>::value, ""); // inferred, like forward
+    auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::identity{});
+    static_assert(is_same<decltype(it_tr_cid)::reference, bool>::value, ""); // special handling by
+                                                                             // transform_iterator_reference
     static_assert(is_same<decltype(it_tr_cid)::value_type, bool>::value, "");
     (void) it_tr_cid;
   }
@@ -234,24 +248,27 @@ void TestTransformIteratorReferenceAndValueType()
     static_assert(is_same<decltype(it_tr_tid)::value_type, bool>::value, "");
     (void) it_tr_tid;
 
-    auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::__identity{});
-    static_assert(is_same<decltype(it_tr_cid)::reference, bool&&>::value, ""); // inferred, like forward
+    auto it_tr_cid = thrust::make_transform_iterator(it, cuda::std::identity{});
+    static_assert(is_same<decltype(it_tr_cid)::reference, bool>::value, ""); // special handling by
+                                                                             // transform_iterator_reference
     static_assert(is_same<decltype(it_tr_cid)::value_type, bool>::value, "");
     (void) it_tr_cid;
   }
+  _CCCL_SUPPRESS_DEPRECATED_POP
 }
 DECLARE_UNITTEST(TestTransformIteratorReferenceAndValueType);
 
 void TestTransformIteratorIdentity()
 {
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
   thrust::device_vector<int> v(3, 42);
 
   ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), thrust::identity<int>{}), 42);
-  // FIXME(bgruber): fix transform_iterator to get these tests compiling:
-  // ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), thrust::identity<>{}), 42);
-  // ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), cuda::std::identity{}), 42);
-  // using namespace thrust::placeholders;
-  // ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), _1), 42);
+  ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), thrust::identity<>{}), 42);
+  ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), cuda::std::identity{}), 42);
+  using namespace thrust::placeholders;
+  ASSERT_EQUAL(*thrust::make_transform_iterator(v.begin(), _1), 42);
+  _CCCL_SUPPRESS_DEPRECATED_POP
 }
 
 DECLARE_UNITTEST(TestTransformIteratorIdentity);
diff --git a/thrust/testing/transform_scan.cu b/thrust/testing/transform_scan.cu
index 66f005a5b93..51971125324 100644
--- a/thrust/testing/transform_scan.cu
+++ b/thrust/testing/transform_scan.cu
@@ -362,20 +362,20 @@ void TestValueCategoryDeduction()
   vec.assign((T*) a_h, a_h + 10);
 
   thrust::transform_inclusive_scan(
-    thrust::device, vec.cbegin(), vec.cend(), vec.begin(), thrust::identity<>{}, thrust::maximum<>{});
+    thrust::device, vec.cbegin(), vec.cend(), vec.begin(), ::cuda::std::identity{}, thrust::maximum<>{});
 
   ASSERT_EQUAL((thrust::device_vector<T>{5, 5, 5, 8, 8, 8, 8, 8, 8, 9}), vec);
 
   vec.assign((T*) a_h, a_h + 10);
 
   thrust::transform_inclusive_scan(
-    thrust::device, vec.cbegin(), vec.cend(), vec.begin(), thrust::identity<>{}, T{}, thrust::maximum<>{});
+    thrust::device, vec.cbegin(), vec.cend(), vec.begin(), ::cuda::std::identity{}, T{}, thrust::maximum<>{});
 
   ASSERT_EQUAL((thrust::device_vector<T>{5, 5, 5, 8, 8, 8, 8, 8, 8, 9}), vec);
 
   vec.assign((T*) a_h, a_h + 10);
   thrust::transform_exclusive_scan(
-    thrust::device, vec.cbegin(), vec.cend(), vec.begin(), thrust::identity<>{}, T{}, thrust::maximum<>{});
+    thrust::device, vec.cbegin(), vec.cend(), vec.begin(), ::cuda::std::identity{}, T{}, thrust::maximum<>{});
 
   ASSERT_EQUAL((thrust::device_vector<T>{0, 5, 5, 5, 8, 8, 8, 8, 8, 8}), vec);
 }
diff --git a/thrust/testing/type_traits.cu b/thrust/testing/type_traits.cu
index db438f05d5e..bbb1989d8d1 100644
--- a/thrust/testing/type_traits.cu
+++ b/thrust/testing/type_traits.cu
@@ -9,6 +9,8 @@
 #include <thrust/tuple.h>
 #include <thrust/type_traits/is_contiguous_iterator.h>
 
+#include <cuda/__cccl_config>
+
 #if _CCCL_COMPILER(GCC, >=, 7)
 // This header pulls in an unsuppressable warning on GCC 6
 #  include <cuda/std/complex>
@@ -36,14 +38,23 @@ void TestIsContiguousIterator()
 
   using HostIteratorTuple = thrust::tuple<HostVector::iterator, HostVector::iterator>;
 
-  using ConstantIterator  = thrust::constant_iterator<int>;
-  using CountingIterator  = thrust::counting_iterator<int>;
-  using TransformIterator = thrust::transform_iterator<thrust::identity<int>, HostVector::iterator>;
-  using ZipIterator       = thrust::zip_iterator<HostIteratorTuple>;
+  using ConstantIterator = thrust::constant_iterator<int>;
+  using CountingIterator = thrust::counting_iterator<int>;
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
+  using TransformIterator1 = thrust::transform_iterator<thrust::identity<int>, HostVector::iterator>;
+  _CCCL_SUPPRESS_DEPRECATED_POP
+  using TransformIterator2 = thrust::transform_iterator<cuda::std::identity, HostVector::iterator>;
+  using ZipIterator        = thrust::zip_iterator<HostIteratorTuple>;
 
   ASSERT_EQUAL((bool) thrust::is_contiguous_iterator<ConstantIterator>::value, false);
   ASSERT_EQUAL((bool) thrust::is_contiguous_iterator<CountingIterator>::value, false);
-  ASSERT_EQUAL((bool) thrust::is_contiguous_iterator<TransformIterator>::value, false);
+#if !_CCCL_COMPILER(NVHPC)
+  // thrust::identity creates a deprecated warning that could not be worked around
+  _CCCL_SUPPRESS_DEPRECATED_PUSH
+  ASSERT_EQUAL((bool) thrust::is_contiguous_iterator<TransformIterator1>::value, false);
+  _CCCL_SUPPRESS_DEPRECATED_POP
+#endif // !_CCCL_COMPILER(NVHPC)
+  ASSERT_EQUAL((bool) thrust::is_contiguous_iterator<TransformIterator2>::value, false);
   ASSERT_EQUAL((bool) thrust::is_contiguous_iterator<ZipIterator>::value, false);
 }
 DECLARE_UNITTEST(TestIsContiguousIterator);
diff --git a/thrust/thrust/detail/functional/actor.h b/thrust/thrust/detail/functional/actor.h
index 2de51b62285..f8eaea8d9fb 100644
--- a/thrust/thrust/detail/functional/actor.h
+++ b/thrust/thrust/detail/functional/actor.h
@@ -35,7 +35,6 @@
 #  pragma system_header
 #endif // no system header
 #include <thrust/detail/type_deduction.h>
-#include <thrust/detail/type_traits/result_of_adaptable_function.h>
 #include <thrust/tuple.h>
 
 #include <cuda/std/type_traits>
@@ -211,11 +210,5 @@ _CCCL_HOST_DEVICE auto compose(Eval e, const SubExpr1& subexpr1, const SubExpr2&
     {{::cuda::std::move(e)}, make_actor(subexpr1), make_actor(subexpr2)}};
 }
 } // namespace functional
-
-template <typename Eval, typename... Args>
-struct result_of_adaptable_function<functional::actor<Eval>(Args...)>
-{
-  using type = decltype(::cuda::std::declval<functional::actor<Eval>>()(::cuda::std::declval<Args>()...));
-};
 } // namespace detail
 THRUST_NAMESPACE_END
diff --git a/thrust/thrust/detail/functional/operators.h b/thrust/thrust/detail/functional/operators.h
index f4f89f5f7d7..dd7b8d979de 100644
--- a/thrust/thrust/detail/functional/operators.h
+++ b/thrust/thrust/detail/functional/operators.h
@@ -34,7 +34,6 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
-#include <thrust/detail/type_traits/result_of_adaptable_function.h>
 #include <thrust/functional.h>
 #include <thrust/tuple.h>
 
diff --git a/thrust/thrust/detail/type_traits/result_of_adaptable_function.h b/thrust/thrust/detail/type_traits/result_of_adaptable_function.h
deleted file mode 100644
index 44138e5f909..00000000000
--- a/thrust/thrust/detail/type_traits/result_of_adaptable_function.h
+++ /dev/null
@@ -1,67 +0,0 @@
-/*
- *  Copyright 2008-2013 NVIDIA Corporation
- *
- *  Licensed under the Apache License, Version 2.0 (the "License");
- *  you may not use this file except in compliance with the License.
- *  You may obtain a copy of the License at
- *
- *      http://www.apache.org/licenses/LICENSE-2.0
- *
- *  Unless required by applicable law or agreed to in writing, software
- *  distributed under the License is distributed on an "AS IS" BASIS,
- *  WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
- *  See the License for the specific language governing permissions and
- *  limitations under the License.
- */
-
-#pragma once
-
-#include <thrust/detail/config.h>
-
-#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC)
-#  pragma GCC system_header
-#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG)
-#  pragma clang system_header
-#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
-#  pragma system_header
-#endif // no system header
-#include <thrust/detail/type_traits.h>
-
-#include <cuda/std/__type_traits/void_t.h>
-
-#include <type_traits>
-
-THRUST_NAMESPACE_BEGIN
-namespace detail
-{
-// Sets `type` to the result of the specified Signature invocation. If the callable defines a `result_type` alias
-// member, that type is used instead. Use invoke_result / result_of when FuncType::result_type is not defined.
-template <typename Signature, typename Enable = void>
-struct result_of_adaptable_function
-{
-private:
-  template <typename Sig>
-  struct impl;
-
-  template <typename F, typename... Args>
-  struct impl<F(Args...)>
-  {
-    using type = invoke_result_t<F, Args...>;
-  };
-
-public:
-  using type = typename impl<Signature>::type;
-};
-
-// TODO(bgruber): remove this specialization eventually
-// specialization for invocations which define result_type
-_CCCL_SUPPRESS_DEPRECATED_PUSH
-template <typename Functor, typename... ArgTypes>
-struct result_of_adaptable_function<Functor(ArgTypes...), ::cuda::std::void_t<typename Functor::result_type>>
-{
-  using type = typename Functor::result_type;
-};
-_CCCL_SUPPRESS_DEPRECATED_POP
-
-} // namespace detail
-THRUST_NAMESPACE_END
diff --git a/thrust/thrust/execution_policy.h b/thrust/thrust/execution_policy.h
index 8f733215e9a..17713553ab0 100644
--- a/thrust/thrust/execution_policy.h
+++ b/thrust/thrust/execution_policy.h
@@ -115,7 +115,7 @@ using device_t = thrust::system::__THRUST_DEVICE_SYSTEM_NAMESPACE::detail::par_t
  *    thrust::for_each(exec, data, data + 4, ignore_argument());
  *
  *    // can't dispatch thrust::transform because no overload exists for my_policy:
- *    //thrust::transform(exec, data, data, + 4, data, thrust::identity<int>()); // error!
+ *    //thrust::transform(exec, data, data, + 4, data, ::cuda::std::identity{}); // error!
  *
  *    return 0;
  *  }
@@ -175,7 +175,7 @@ struct execution_policy : thrust::detail::execution_policy_base<DerivedPolicy>
  *    thrust::for_each(exec, data, data + 4, ignore_argument());
  *
  *    // dispatch thrust::transform whose behavior our policy inherits
- *    thrust::transform(exec, data, data, + 4, data, thrust::identity<int>());
+ *    thrust::transform(exec, data, data, + 4, data, ::cuda::std::identity{});
  *
  *    return 0;
  *  }
@@ -234,7 +234,7 @@ struct host_execution_policy : thrust::system::__THRUST_HOST_SYSTEM_NAMESPACE::e
  *    thrust::for_each(exec, data, data + 4, ignore_argument());
  *
  *    // dispatch thrust::transform whose behavior our policy inherits
- *    thrust::transform(exec, data, data, + 4, data, thrust::identity<int>());
+ *    thrust::transform(exec, data, data, + 4, data, ::cuda::std::identity{});
  *
  *    return 0;
  *  }
diff --git a/thrust/thrust/functional.h b/thrust/thrust/functional.h
index 816c8744b89..817f2bcb443 100644
--- a/thrust/thrust/functional.h
+++ b/thrust/thrust/functional.h
@@ -839,7 +839,7 @@ struct bit_xor : public ::cuda::std::bit_xor<T>
  */
 // TODO(bgruber): this version can also act as a functor casting to T making it not equivalent to ::cuda::std::identity
 template <typename T = void>
-struct identity
+struct CCCL_DEPRECATED_BECAUSE("use cuda::std::identity instead") identity
 {
   using argument_type _LIBCUDACXX_DEPRECATED_IN_CXX11 = T;
   using result_type _LIBCUDACXX_DEPRECATED_IN_CXX11   = T;
@@ -866,9 +866,11 @@ struct identity
   }
 };
 
+_CCCL_SUPPRESS_DEPRECATED_PUSH
 template <>
-struct identity<void> : ::cuda::std::__identity
+struct CCCL_DEPRECATED_BECAUSE("use cuda::std::identity instead") identity<void> : ::cuda::std::__identity
 {};
+_CCCL_SUPPRESS_DEPRECATED_POP
 
 /*! \p maximum is a function object that takes two arguments and returns the greater
  *  of the two. Specifically, it is an Adaptable Binary Function. If \c f is an
diff --git a/thrust/thrust/iterator/counting_iterator.h b/thrust/thrust/iterator/counting_iterator.h
index bb98bc3b4d5..689b81b2530 100644
--- a/thrust/thrust/iterator/counting_iterator.h
+++ b/thrust/thrust/iterator/counting_iterator.h
@@ -122,7 +122,7 @@ THRUST_NAMESPACE_BEGIN
  *                                               thrust::make_counting_iterator(8),
  *                                               stencil.begin(),
  *                                               indices.begin(),
- *                                               thrust::identity<int>());
+ *                                               ::cuda::std::identity{});
  *   // indices now contains [1,2,5,7]
  *
  *   return 0;
diff --git a/thrust/thrust/iterator/detail/transform_iterator.inl b/thrust/thrust/iterator/detail/transform_iterator.inl
index 68145e65c50..f01df472f2d 100644
--- a/thrust/thrust/iterator/detail/transform_iterator.inl
+++ b/thrust/thrust/iterator/detail/transform_iterator.inl
@@ -25,10 +25,12 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
-#include <thrust/detail/type_traits/result_of_adaptable_function.h>
+#include <thrust/detail/functional/actor.h>
+#include <thrust/functional.h>
 #include <thrust/iterator/iterator_adaptor.h>
 #include <thrust/iterator/iterator_traits.h>
 
+#include <cuda/std/functional>
 #include <cuda/std/type_traits>
 
 THRUST_NAMESPACE_BEGIN
@@ -39,19 +41,45 @@ class transform_iterator;
 namespace detail
 {
 
+template <class UnaryFunc, class Iterator>
+struct transform_iterator_reference
+{
+  // by default, dereferencing the iterator yields the same as the function.
+  using type = decltype(::cuda::std::declval<UnaryFunc>()(::cuda::std::declval<iterator_value_t<Iterator>>()));
+};
+
+// for certain function objects, we need to tweak the reference type. Notably, identity functions must decay to values.
+// See the implementation of transform_iterator<...>::dereference() for several comments on why this is necessary.
+_CCCL_SUPPRESS_DEPRECATED_PUSH
+template <typename T, class Iterator>
+struct transform_iterator_reference<identity<T>, Iterator>
+{
+  using type = T;
+};
+template <class Iterator>
+struct transform_iterator_reference<identity<>, Iterator>
+{
+  using type = iterator_value_t<Iterator>;
+};
+_CCCL_SUPPRESS_DEPRECATED_POP
+template <class Iterator>
+struct transform_iterator_reference<::cuda::std::identity, Iterator>
+{
+  using type = iterator_value_t<Iterator>;
+};
+template <typename Eval, class Iterator>
+struct transform_iterator_reference<functional::actor<Eval>, Iterator>
+{
+  using type = ::cuda::std::remove_reference_t<decltype(::cuda::std::declval<functional::actor<Eval>>()(
+    ::cuda::std::declval<iterator_value_t<Iterator>>()))>;
+};
+
 // Type function to compute the iterator_adaptor instantiation to be used for transform_iterator
 template <class UnaryFunc, class Iterator, class Reference, class Value>
 struct make_transform_iterator_base
 {
 private:
-  // FIXME(bgruber): the next line should be correct, but thrust::identity<T> lies and advertises a ::return_type of T,
-  // while its operator() returns const T& (which __invoke_of correctly detects), which causes transform_iterator to
-  // crash (or cause UB) during dereferencing. Check the test `thrust.test.dereference` for the OMP and TBB backends.
-  // using wrapped_func_ret_t = ::cuda::std::__invoke_of<UnaryFunc, iterator_value_t<Iterator>>;
-  using wrapped_func_ret_t = result_of_adaptable_function<UnaryFunc(iterator_value_t<Iterator>)>;
-
-  // By default, dereferencing the iterator yields the same as the function.
-  using reference  = typename ia_dflt_help<Reference, wrapped_func_ret_t>::type;
+  using reference  = typename ia_dflt_help<Reference, transform_iterator_reference<UnaryFunc, Iterator>>::type;
   using value_type = typename ia_dflt_help<Value, ::cuda::std::remove_cvref<reference>>::type;
 
 public:
diff --git a/thrust/thrust/logical.h b/thrust/thrust/logical.h
index 747af38554e..a2be3e6087b 100644
--- a/thrust/thrust/logical.h
+++ b/thrust/thrust/logical.h
@@ -65,11 +65,11 @@ THRUST_NAMESPACE_BEGIN
  *  ...
  *  bool A[3] = {true, true, false};
  *
- *  thrust::all_of(thrust::host, A, A + 2, thrust::identity<bool>()); // returns true
- *  thrust::all_of(thrust::host, A, A + 3, thrust::identity<bool>()); // returns false
+ *  thrust::all_of(thrust::host, A, A + 2, ::cuda::std::identity{}); // returns true
+ *  thrust::all_of(thrust::host, A, A + 3, ::cuda::std::identity{}); // returns false
  *
  *  // empty range
- *  thrust::all_of(thrust::host, A, A, thrust::identity<bool>()); // returns false
+ *  thrust::all_of(thrust::host, A, A, ::cuda::std::identity{}); // returns false
  *
  *  \endcode
  *
@@ -104,11 +104,11 @@ all_of(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
  *  ...
  *  bool A[3] = {true, true, false};
  *
- *  thrust::all_of(A, A + 2, thrust::identity<bool>()); // returns true
- *  thrust::all_of(A, A + 3, thrust::identity<bool>()); // returns false
+ *  thrust::all_of(A, A + 2, ::cuda::std::identity{}); // returns true
+ *  thrust::all_of(A, A + 3, ::cuda::std::identity{}); // returns false
  *
  *  // empty range
- *  thrust::all_of(A, A, thrust::identity<bool>()); // returns false
+ *  thrust::all_of(A, A, ::cuda::std::identity{}); // returns false
  *
  *  \endcode
  *
@@ -144,13 +144,13 @@ bool all_of(InputIterator first, InputIterator last, Predicate pred);
  *  ...
  *  bool A[3] = {true, true, false};
  *
- *  thrust::any_of(thrust::host, A, A + 2, thrust::identity<bool>()); // returns true
- *  thrust::any_of(thrust::host, A, A + 3, thrust::identity<bool>()); // returns true
+ *  thrust::any_of(thrust::host, A, A + 2, ::cuda::std::identity{}); // returns true
+ *  thrust::any_of(thrust::host, A, A + 3, ::cuda::std::identity{}); // returns true
  *
- *  thrust::any_of(thrust::host, A + 2, A + 3, thrust::identity<bool>()); // returns false
+ *  thrust::any_of(thrust::host, A + 2, A + 3, ::cuda::std::identity{}); // returns false
  *
  *  // empty range
- *  thrust::any_of(thrust::host, A, A, thrust::identity<bool>()); // returns false
+ *  thrust::any_of(thrust::host, A, A, ::cuda::std::identity{}); // returns false
  *  \endcode
  *
  *  \see all_of
@@ -184,13 +184,13 @@ any_of(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
  *  ...
  *  bool A[3] = {true, true, false};
  *
- *  thrust::any_of(A, A + 2, thrust::identity<bool>()); // returns true
- *  thrust::any_of(A, A + 3, thrust::identity<bool>()); // returns true
+ *  thrust::any_of(A, A + 2, ::cuda::std::identity{}); // returns true
+ *  thrust::any_of(A, A + 3, ::cuda::std::identity{}); // returns true
  *
- *  thrust::any_of(A + 2, A + 3, thrust::identity<bool>()); // returns false
+ *  thrust::any_of(A + 2, A + 3, ::cuda::std::identity{}); // returns false
  *
  *  // empty range
- *  thrust::any_of(A, A, thrust::identity<bool>()); // returns false
+ *  thrust::any_of(A, A, ::cuda::std::identity{}); // returns false
  *  \endcode
  *
  *  \see all_of
@@ -225,13 +225,13 @@ bool any_of(InputIterator first, InputIterator last, Predicate pred);
  *  ...
  *  bool A[3] = {true, true, false};
  *
- *  thrust::none_of(thrust::host, A, A + 2, thrust::identity<bool>()); // returns false
- *  thrust::none_of(thrust::host, A, A + 3, thrust::identity<bool>()); // returns false
+ *  thrust::none_of(thrust::host, A, A + 2, ::cuda::std::identity{}); // returns false
+ *  thrust::none_of(thrust::host, A, A + 3, ::cuda::std::identity{}); // returns false
  *
- *  thrust::none_of(thrust::host, A + 2, A + 3, thrust::identity<bool>()); // returns true
+ *  thrust::none_of(thrust::host, A + 2, A + 3, ::cuda::std::identity{}); // returns true
  *
  *  // empty range
- *  thrust::none_of(thrust::host, A, A, thrust::identity<bool>()); // returns true
+ *  thrust::none_of(thrust::host, A, A, ::cuda::std::identity{}); // returns true
  *  \endcode
  *
  *  \see all_of
@@ -265,13 +265,13 @@ none_of(const thrust::detail::execution_policy_base<DerivedPolicy>& exec,
  *  ...
  *  bool A[3] = {true, true, false};
  *
- *  thrust::none_of(A, A + 2, thrust::identity<bool>()); // returns false
- *  thrust::none_of(A, A + 3, thrust::identity<bool>()); // returns false
+ *  thrust::none_of(A, A + 2, ::cuda::std::identity{}); // returns false
+ *  thrust::none_of(A, A + 3, ::cuda::std::identity{}); // returns false
  *
- *  thrust::none_of(A + 2, A + 3, thrust::identity<bool>()); // returns true
+ *  thrust::none_of(A + 2, A + 3, ::cuda::std::identity{}); // returns true
  *
  *  // empty range
- *  thrust::none_of(A, A, thrust::identity<bool>()); // returns true
+ *  thrust::none_of(A, A, ::cuda::std::identity{}); // returns true
  *  \endcode
  *
  *  \see all_of
diff --git a/thrust/thrust/partition.h b/thrust/thrust/partition.h
index a9b3a4d1989..11515010796 100644
--- a/thrust/thrust/partition.h
+++ b/thrust/thrust/partition.h
@@ -485,7 +485,7 @@ thrust::pair<OutputIterator1, OutputIterator2> partition_copy(
  *  const int N = sizeof(A)/sizeof(int);
  *  int *evens = result;
  *  int *odds  = result + 5;
- *  thrust::stable_partition_copy(thrust::host, A, A + N, S, evens, odds, thrust::identity<int>());
+ *  thrust::stable_partition_copy(thrust::host, A, A + N, S, evens, odds, ::cuda::std::identity{});
  *  // A remains {1, 2, 3, 4, 5, 6, 7, 8, 9, 10}
  *  // S remains {0, 1, 0, 1, 0, 1, 0, 1, 0,  1}
  *  // result is now {2, 4, 6, 8, 10, 1, 3, 5, 7, 9}
@@ -560,7 +560,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> partition_copy(
  *  const int N = sizeof(A)/sizeof(int);
  *  int *evens = result;
  *  int *odds  = result + 5;
- *  thrust::stable_partition_copy(A, A + N, S, evens, odds, thrust::identity<int>());
+ *  thrust::stable_partition_copy(A, A + N, S, evens, odds, ::cuda::std::identity{});
  *  // A remains {1, 2, 3, 4, 5, 6, 7, 8, 9, 10}
  *  // S remains {0, 1, 0, 1, 0, 1, 0, 1, 0,  1}
  *  // result is now {2, 4, 6, 8, 10, 1, 3, 5, 7, 9}
@@ -1052,7 +1052,7 @@ thrust::pair<OutputIterator1, OutputIterator2> stable_partition_copy(
  *  const int N = sizeof(A)/sizeof(int);
  *  int *evens = result;
  *  int *odds  = result + 5;
- *  thrust::stable_partition_copy(thrust::host, A, A + N, S, evens, odds, thrust::identity<int>());
+ *  thrust::stable_partition_copy(thrust::host, A, A + N, S, evens, odds, ::cuda::std::identity{});
  *  // A remains {1, 2, 3, 4, 5, 6, 7, 8, 9, 10}
  *  // S remains {0, 1, 0, 1, 0, 1, 0, 1, 0,  1}
  *  // result is now {2, 4, 6, 8, 10, 1, 3, 5, 7, 9}
@@ -1129,7 +1129,7 @@ _CCCL_HOST_DEVICE thrust::pair<OutputIterator1, OutputIterator2> stable_partitio
  *  const int N = sizeof(A)/sizeof(int);
  *  int *evens = result;
  *  int *odds  = result + 5;
- *  thrust::stable_partition_copy(A, A + N, S, evens, odds, thrust::identity<int>());
+ *  thrust::stable_partition_copy(A, A + N, S, evens, odds, ::cuda::std::identity{});
  *  // A remains {1, 2, 3, 4, 5, 6, 7, 8, 9, 10}
  *  // S remains {0, 1, 0, 1, 0, 1, 0, 1, 0,  1}
  *  // result is now {2, 4, 6, 8, 10, 1, 3, 5, 7, 9}
diff --git a/thrust/thrust/remove.h b/thrust/thrust/remove.h
index 80dfa786dcc..260f8a32947 100644
--- a/thrust/thrust/remove.h
+++ b/thrust/thrust/remove.h
@@ -565,7 +565,7 @@ OutputIterator remove_copy_if(InputIterator first, InputIterator last, OutputIte
  *  int A[N] = {1, 4, 2, 8, 5, 7};
  *  int S[N] = {0, 1, 1, 1, 0, 0};
  *
- *  int *new_end = thrust::remove_if(thrust::host, A, A + N, S, thrust::identity<int>());
+ *  int *new_end = thrust::remove_if(thrust::host, A, A + N, S, ::cuda::std::identity{});
  *  // The first three values of A are now {1, 5, 7}
  *  // Values beyond new_end are unspecified
  *  \endcode
@@ -623,7 +623,7 @@ _CCCL_HOST_DEVICE ForwardIterator remove_if(
  *  int A[N] = {1, 4, 2, 8, 5, 7};
  *  int S[N] = {0, 1, 1, 1, 0, 0};
  *
- *  int *new_end = thrust::remove_if(A, A + N, S, thrust::identity<int>());
+ *  int *new_end = thrust::remove_if(A, A + N, S, ::cuda::std::identity{});
  *  // The first three values of A are now {1, 5, 7}
  *  // Values beyond new_end are unspecified
  *  \endcode
@@ -683,7 +683,7 @@ ForwardIterator remove_if(ForwardIterator first, ForwardIterator last, InputIter
  *  int V[N] = {-2, 0, -1, 0, 1, 2};
  *  int S[N] = { 1, 1,  0, 1, 0, 1};
  *  int result[2];
- *  thrust::remove_copy_if(thrust::host, V, V + N, S, result, thrust::identity<int>());
+ *  thrust::remove_copy_if(thrust::host, V, V + N, S, result, ::cuda::std::identity{});
  *  // V remains {-2, 0, -1, 0, 1, 2}
  *  // result is now {-1, 1}
  *  \endcode
@@ -745,7 +745,7 @@ _CCCL_HOST_DEVICE OutputIterator remove_copy_if(
  *  int V[N] = {-2, 0, -1, 0, 1, 2};
  *  int S[N] = { 1, 1,  0, 1, 0, 1};
  *  int result[2];
- *  thrust::remove_copy_if(V, V + N, S, result, thrust::identity<int>());
+ *  thrust::remove_copy_if(V, V + N, S, result, ::cuda::std::identity{});
  *  // V remains {-2, 0, -1, 0, 1, 2}
  *  // result is now {-1, 1}
  *  \endcode
diff --git a/thrust/thrust/system/cuda/detail/async/copy.h b/thrust/thrust/system/cuda/detail/async/copy.h
index b4ec3e754c0..e98bb4a49a2 100644
--- a/thrust/thrust/system/cuda/detail/async/copy.h
+++ b/thrust/thrust/system/cuda/detail/async/copy.h
@@ -129,9 +129,7 @@ auto async_copy_n(thrust::cuda::execution_policy<FromPolicy>& from_exec,
                                       decltype(is_device_to_device_copy(from_exec, to_exec))>::value,
                           unique_eager_event>::type
 {
-  using T = typename iterator_traits<ForwardIt>::value_type;
-
-  return async_transform_n(select_device_system(from_exec, to_exec), first, n, output, thrust::identity<T>());
+  return async_transform_n(select_device_system(from_exec, to_exec), first, n, output, ::cuda::std::identity{});
 }
 
 template <typename OutputIt>
diff --git a/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h b/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
index c3ab229e413..0cb64431956 100644
--- a/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
+++ b/thrust/thrust/system/cuda/detail/internal/copy_device_to_device.h
@@ -81,8 +81,7 @@ template <class Derived, class InputIt, class OutputIt>
 OutputIt THRUST_RUNTIME_FUNCTION device_to_device(
   execution_policy<Derived>& policy, InputIt first, InputIt last, OutputIt result, thrust::detail::false_type)
 {
-  using InputTy = typename thrust::iterator_traits<InputIt>::value_type;
-  return cuda_cub::transform(policy, first, last, result, thrust::identity<InputTy>());
+  return cuda_cub::transform(policy, first, last, result, ::cuda::std::identity{});
 }
 
 template <class Derived, class InputIt, class OutputIt>
diff --git a/thrust/thrust/system/cuda/detail/scatter.h b/thrust/thrust/system/cuda/detail/scatter.h
index 23976ed57b4..17a61c70f77 100644
--- a/thrust/thrust/system/cuda/detail/scatter.h
+++ b/thrust/thrust/system/cuda/detail/scatter.h
@@ -40,7 +40,7 @@
 #  include <thrust/iterator/permutation_iterator.h>
 #  include <thrust/system/cuda/detail/transform.h>
 
-#  include <cuda/std/__functional/identity.h>
+#  include <cuda/std/functional>
 
 THRUST_NAMESPACE_BEGIN
 namespace cuda_cub
diff --git a/thrust/thrust/system/detail/generic/copy.inl b/thrust/thrust/system/detail/generic/copy.inl
index 8a022a16c68..f4ce9b0c92b 100644
--- a/thrust/thrust/system/detail/generic/copy.inl
+++ b/thrust/thrust/system/detail/generic/copy.inl
@@ -46,16 +46,14 @@ template <typename DerivedPolicy, typename InputIterator, typename OutputIterato
 _CCCL_HOST_DEVICE OutputIterator
 copy(thrust::execution_policy<DerivedPolicy>& exec, InputIterator first, InputIterator last, OutputIterator result)
 {
-  using T = typename thrust::iterator_value<InputIterator>::type;
-  return thrust::transform(exec, first, last, result, thrust::identity<T>());
+  return thrust::transform(exec, first, last, result, ::cuda::std::identity{});
 } // end copy()
 
 template <typename DerivedPolicy, typename InputIterator, typename Size, typename OutputIterator>
 _CCCL_HOST_DEVICE OutputIterator
 copy_n(thrust::execution_policy<DerivedPolicy>& exec, InputIterator first, Size n, OutputIterator result)
 {
-  using value_type = typename thrust::iterator_value<InputIterator>::type;
-  using xfrm_type  = thrust::identity<value_type>;
+  using xfrm_type = ::cuda::std::identity;
 
   using functor_type = thrust::detail::unary_transform_functor<xfrm_type>;
 
diff --git a/thrust/thrust/system/detail/generic/copy_if.inl b/thrust/thrust/system/detail/generic/copy_if.inl
index 6ccc52732c1..2264a0f5544 100644
--- a/thrust/thrust/system/detail/generic/copy_if.inl
+++ b/thrust/thrust/system/detail/generic/copy_if.inl
@@ -83,8 +83,7 @@ _CCCL_HOST_DEVICE OutputIterator copy_if(
     thrust::plus<IndexType>());
 
   // scatter the true elements
-  thrust::scatter_if(
-    exec, first, last, scatter_indices.begin(), predicates.begin(), result, thrust::identity<IndexType>());
+  thrust::scatter_if(exec, first, last, scatter_indices.begin(), predicates.begin(), result);
 
   // find the end of the new sequence
   IndexType output_size = scatter_indices[n - 1] + predicates[n - 1];
diff --git a/thrust/thrust/system/detail/generic/gather.inl b/thrust/thrust/system/detail/generic/gather.inl
index 569a7fbe366..c24c9ec52f8 100644
--- a/thrust/thrust/system/detail/generic/gather.inl
+++ b/thrust/thrust/system/detail/generic/gather.inl
@@ -52,7 +52,7 @@ _CCCL_HOST_DEVICE OutputIterator gather(
     thrust::make_permutation_iterator(input_first, map_first),
     thrust::make_permutation_iterator(input_first, map_last),
     result,
-    thrust::identity<typename thrust::iterator_value<RandomAccessIterator>::type>());
+    ::cuda::std::identity{});
 } // end gather()
 
 template <typename DerivedPolicy,
@@ -68,8 +68,7 @@ _CCCL_HOST_DEVICE OutputIterator gather_if(
   RandomAccessIterator input_first,
   OutputIterator result)
 {
-  using StencilType = typename thrust::iterator_value<InputIterator2>::type;
-  return thrust::gather_if(exec, map_first, map_last, stencil, input_first, result, thrust::identity<StencilType>());
+  return thrust::gather_if(exec, map_first, map_last, stencil, input_first, result, ::cuda::std::identity{});
 } // end gather_if()
 
 template <typename DerivedPolicy,
@@ -87,14 +86,13 @@ _CCCL_HOST_DEVICE OutputIterator gather_if(
   OutputIterator result,
   Predicate pred)
 {
-  using InputType = typename thrust::iterator_value<RandomAccessIterator>::type;
   return thrust::transform_if(
     exec,
     thrust::make_permutation_iterator(input_first, map_first),
     thrust::make_permutation_iterator(input_first, map_last),
     stencil,
     result,
-    thrust::identity<InputType>(),
+    ::cuda::std::identity{},
     pred);
 } // end gather_if()
 
diff --git a/thrust/thrust/system/detail/generic/scatter.inl b/thrust/thrust/system/detail/generic/scatter.inl
index 20298f0b7b2..72216df1275 100644
--- a/thrust/thrust/system/detail/generic/scatter.inl
+++ b/thrust/thrust/system/detail/generic/scatter.inl
@@ -25,11 +25,11 @@
 #elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC)
 #  pragma system_header
 #endif // no system header
+#include <thrust/copy.h>
 #include <thrust/functional.h>
 #include <thrust/iterator/iterator_traits.h>
 #include <thrust/iterator/permutation_iterator.h>
 #include <thrust/system/detail/generic/scatter.h>
-#include <thrust/transform.h>
 
 THRUST_NAMESPACE_BEGIN
 namespace system
@@ -47,12 +47,7 @@ scatter(thrust::execution_policy<DerivedPolicy>& exec,
         InputIterator2 map,
         RandomAccessIterator output)
 {
-  thrust::transform(
-    exec,
-    first,
-    last,
-    thrust::make_permutation_iterator(output, map),
-    thrust::identity<typename thrust::iterator_value<InputIterator1>::type>());
+  thrust::copy(exec, first, last, thrust::make_permutation_iterator(output, map));
 } // end scatter()
 
 template <typename DerivedPolicy,
@@ -68,9 +63,7 @@ _CCCL_HOST_DEVICE void scatter_if(
   InputIterator3 stencil,
   RandomAccessIterator output)
 {
-  // default predicate is identity
-  using StencilType = typename thrust::iterator_value<InputIterator3>::type;
-  thrust::scatter_if(exec, first, last, map, stencil, output, thrust::identity<StencilType>());
+  thrust::scatter_if(exec, first, last, map, stencil, output, ::cuda::std::identity{});
 } // end scatter_if()
 
 template <typename DerivedPolicy,
@@ -88,9 +81,8 @@ _CCCL_HOST_DEVICE void scatter_if(
   RandomAccessIterator output,
   Predicate pred)
 {
-  using InputType = typename thrust::iterator_value<InputIterator1>::type;
   thrust::transform_if(
-    exec, first, last, stencil, thrust::make_permutation_iterator(output, map), thrust::identity<InputType>(), pred);
+    exec, first, last, stencil, thrust::make_permutation_iterator(output, map), ::cuda::std::identity{}, pred);
 } // end scatter_if()
 
 } // end namespace generic
diff --git a/thrust/thrust/transform.h b/thrust/thrust/transform.h
index 9c96f9ca6f5..cb5170c3407 100644
--- a/thrust/thrust/transform.h
+++ b/thrust/thrust/transform.h
@@ -327,7 +327,6 @@ transform(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, Ou
  *  };
  *
  *  thrust::negate<int> op;
- *  thrust::identity<int> identity;
  *
  *  // negate odd elements
  *  thrust::transform_if(thrust::host, data, data + 10, data, op, is_odd()); // in-place transformation
@@ -402,7 +401,6 @@ _CCCL_HOST_DEVICE ForwardIterator transform_if(
  *  };
  *
  *  thrust::negate<int> op;
- *  thrust::identity<int> identity;
  *
  *  // negate odd elements
  *  thrust::transform_if(data, data + 10, data, op, is_odd()); // in-place transformation
@@ -472,7 +470,7 @@ transform_if(InputIterator first, InputIterator last, ForwardIterator result, Un
  *  int stencil[10] = { 1, 0, 1,  0, 1, 0, 1,  0, 1, 0};
  *
  *  thrust::negate<int> op;
- *  thrust::identity<int> identity;
+ *  ::cuda::std::identity identity;
  *
  *  thrust::transform_if(thrust::host, data, data + 10, stencil, data, op, identity); // in-place transformation
  *
@@ -545,7 +543,7 @@ _CCCL_HOST_DEVICE ForwardIterator transform_if(
  *  int stencil[10] = { 1, 0, 1,  0, 1, 0, 1,  0, 1, 0};
  *
  *  thrust::negate<int> op;
- *  thrust::identity<int> identity;
+ *  ::cuda::std::identity identity;
  *
  *  thrust::transform_if(data, data + 10, stencil, data, op, identity); // in-place transformation
  *
@@ -628,7 +626,7 @@ ForwardIterator transform_if(
  *  int output[6];
  *
  *  thrust::plus<int> op;
- *  thrust::identity<int> identity;
+ *  ::cuda::std::identity identity;
  *
  *  thrust::transform_if(thrust::host, input1, input1 + 6, input2, stencil, output, op, identity);
  *
@@ -709,7 +707,7 @@ _CCCL_HOST_DEVICE ForwardIterator transform_if(
  *  int output[6];
  *
  *  thrust::plus<int> op;
- *  thrust::identity<int> identity;
+ *  ::cuda::std::identity identity;
  *
  *  thrust::transform_if(input1, input1 + 6, input2, stencil, output, op, identity);
  *