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 Hash = cuda::std::identity, typename KeyEqual = thrust::equal_to, 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::iterator; - IndexIterator indices_end = thrust::copy_if(first, last, stencil.begin(), indices.begin(), thrust::identity()); + 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..0fc73ea84a2 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()); + thrust::for_each(sys, vec.begin(), vec.end(), ::cuda::std::identity{}); // 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()); + 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()); + thrust::for_each(vec.begin(), vec.end(), ::cuda::std::identity{}); 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()); + 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 35f59b13ffa..aac15bbcb0a 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())); + 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())); + 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())); + 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())); + 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())); + 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())); + 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 v(3, 1); thrust::device_vector result(1); - all_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity(), 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(), 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(), 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(), 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(), 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(), 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()), 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()), 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()), true); - ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, thrust::identity()), true); - ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, thrust::identity()), false); - ASSERT_EQUAL(thrust::all_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, thrust::identity()), 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 v(3, 1); thrust::device_vector result(1); - any_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity(), 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(), 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(), 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(), 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(), 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(), 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()), 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()), 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()), false); - ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, thrust::identity()), true); - ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, thrust::identity()), true); - ASSERT_EQUAL(thrust::any_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, thrust::identity()), 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 v(3, 1); thrust::device_vector result(1); - none_of_kernel<<<1, 1>>>(exec, v.begin(), v.end(), thrust::identity(), 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(), 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(), 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(), 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(), 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(), 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()), 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()), 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()), true); - ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 1, thrust::identity()), false); - ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 0, v.begin() + 2, thrust::identity()), false); - ASSERT_EQUAL(thrust::none_of(thrust::cuda::par.on(s), v.begin() + 1, v.begin() + 2, thrust::identity()), 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())); + 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())); + 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()); + 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()); + 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()); + 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()); + 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 594194cd183..ada3c2dae8c 100644 --- a/thrust/testing/cuda/transform.cu +++ b/thrust/testing/cuda/transform.cu @@ -80,7 +80,7 @@ void TestTransformIfUnaryNoStencilDevice(ExecutionPolicy exec) thrust::device_vector iter_vec(1); transform_if_kernel<<<1, 1>>>( - exec, input.begin(), input.end(), output.begin(), thrust::negate(), thrust::identity(), iter_vec.begin()); + exec, input.begin(), input.end(), output.begin(), thrust::negate(), ::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(), - thrust::identity(), + ::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 identity; + ::cuda::std::identity identity; thrust::device_vector 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(), thrust::plus()); + thrust::transform_inclusive_scan( + input.begin(), input.end(), output.begin(), ::cuda::std::identity{}, thrust::plus()); thrust::inclusive_scan(input.begin(), input.end(), reference.begin(), thrust::plus()); 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 input = unittest::random_integers(100); thrust::device_vector output(input.size(), 0); - simple_copy(thrust::make_transform_iterator(input.begin(), thrust::identity()), - thrust::make_transform_iterator(input.end(), thrust::identity()), + 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 input = unittest::random_integers(100); thrust::device_vector output(input.size(), 0); - simple_copy(thrust::make_transform_iterator(input.begin(), thrust::identity()), - thrust::make_transform_iterator(input.end(), thrust::identity()), + 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 input = unittest::random_integers(100); thrust::device_vector output(input.size(), 0); - simple_copy(thrust::make_transform_iterator(input.begin(), thrust::identity()), - thrust::make_transform_iterator(input.end(), thrust::identity()), + 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 24894fa0026..f343461ccd4 100644 --- a/thrust/testing/find.cu +++ b/thrust/testing/find.cu @@ -128,7 +128,7 @@ void TestFindIfDispatchExplicit() thrust::device_vector vec(1); my_system sys(0); - thrust::find_if(sys, vec.begin(), vec.end(), thrust::identity()); + thrust::find_if(sys, vec.begin(), vec.end(), ::cuda::std::identity{}); ASSERT_EQUAL(true, sys.is_valid()); } @@ -145,7 +145,7 @@ void TestFindIfDispatchImplicit() { thrust::device_vector vec(1); - thrust::find_if(thrust::retag(vec.begin()), thrust::retag(vec.end()), thrust::identity()); + thrust::find_if(thrust::retag(vec.begin()), thrust::retag(vec.end()), ::cuda::std::identity{}); ASSERT_EQUAL(13, vec.front()); } @@ -179,7 +179,7 @@ void TestFindIfNotDispatchExplicit() thrust::device_vector vec(1); my_system sys(0); - thrust::find_if_not(sys, vec.begin(), vec.end(), thrust::identity()); + thrust::find_if_not(sys, vec.begin(), vec.end(), ::cuda::std::identity{}); ASSERT_EQUAL(true, sys.is_valid()); } @@ -196,7 +196,7 @@ void TestFindIfNotDispatchImplicit() { thrust::device_vector vec(1); - thrust::find_if_not(thrust::retag(vec.begin()), thrust::retag(vec.end()), thrust::identity()); + thrust::find_if_not(thrust::retag(vec.begin()), thrust::retag(vec.end()), ::cuda::std::identity{}); ASSERT_EQUAL(13, vec.front()); } diff --git a/thrust/testing/functional.cu b/thrust/testing/functional.cu index 1818084fb34..fc5d680dfea 100644 --- a/thrust/testing/functional.cu +++ b/thrust/testing/functional.cu @@ -193,6 +193,7 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestIdentityFunctional() int i = 42; double d = 3.14; + // TODO // pass through ASSERT_EQUAL(thrust::identity{}(i), 42); ASSERT_EQUAL(thrust::identity{}(d), 3); @@ -228,7 +229,7 @@ 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()); + thrust::transform(input.begin(), input.end(), output.begin(), ::cuda::std::identity{}); ASSERT_EQUAL(input, output); } DECLARE_VECTOR_UNITTEST(TestIdentityFunctionalVector); @@ -308,7 +309,7 @@ THRUST_DISABLE_BROKEN_GCC_VECTORIZER void TestNot1() Vector output(5); - thrust::transform(input.begin(), input.end(), output.begin(), thrust::not_fn(thrust::identity())); + 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..581e78715b8 100644 --- a/thrust/testing/is_partitioned.cu +++ b/thrust/testing/is_partitioned.cu @@ -21,24 +21,24 @@ void TestIsPartitionedSimple() Vector v{1, 1, 1, 0}; // empty partition - ASSERT_EQUAL_QUIET(true, thrust::is_partitioned(v.begin(), v.begin(), thrust::identity())); + 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())); + 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())); + 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())); + 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())); + 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())); + 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()), 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()), 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()), true); - ASSERT_EQUAL(thrust::all_of(v.begin() + 0, v.begin() + 1, thrust::identity()), true); - ASSERT_EQUAL(thrust::all_of(v.begin() + 0, v.begin() + 2, thrust::identity()), false); - ASSERT_EQUAL(thrust::all_of(v.begin() + 1, v.begin() + 2, thrust::identity()), 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()), 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()), 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()), false); - ASSERT_EQUAL(thrust::any_of(v.begin() + 0, v.begin() + 1, thrust::identity()), true); - ASSERT_EQUAL(thrust::any_of(v.begin() + 0, v.begin() + 2, thrust::identity()), true); - ASSERT_EQUAL(thrust::any_of(v.begin() + 1, v.begin() + 2, thrust::identity()), 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()), 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()), 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()), true); - ASSERT_EQUAL(thrust::none_of(v.begin() + 0, v.begin() + 1, thrust::identity()), false); - ASSERT_EQUAL(thrust::none_of(v.begin() + 0, v.begin() + 2, thrust::identity()), false); - ASSERT_EQUAL(thrust::none_of(v.begin() + 1, v.begin() + 2, thrust::identity()), 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..38bd0763217 100644 --- a/thrust/testing/partition.cu +++ b/thrust/testing/partition.cu @@ -146,7 +146,7 @@ void TestStablePartitionStencilSimple() Vector stencil{0, 1, 0, 0, 1}; - Iterator iter = thrust::stable_partition(data.begin(), data.end(), stencil.begin(), thrust::identity()); + Iterator iter = thrust::stable_partition(data.begin(), data.end(), stencil.begin(), ::cuda::std::identity{}); Vector ref{2, 2, 1, 1, 3}; @@ -192,7 +192,7 @@ void TestStablePartitionCopyStencilSimple() Vector false_results(3); thrust::pair ends = thrust::stable_partition_copy( - data.begin(), data.end(), stencil.begin(), true_results.begin(), false_results.begin(), thrust::identity()); + 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..9b6076034f6 100644 --- a/thrust/testing/partition_point.cu +++ b/thrust/testing/partition_point.cu @@ -25,11 +25,11 @@ void TestPartitionPointSimple() Iterator last = v.begin() + 4; Iterator ref = first + 3; - ASSERT_EQUAL_QUIET(ref, thrust::partition_point(first, last, thrust::identity())); + 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())); + 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()); + ::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..097e9ac7373 100644 --- a/thrust/testing/remove.cu +++ b/thrust/testing/remove.cu @@ -195,7 +195,7 @@ void TestRemoveIfStencilSimple() Vector stencil{0, 1, 0, 0, 1}; - typename Vector::iterator end = thrust::remove_if(data.begin(), data.end(), stencil.begin(), thrust::identity()); + 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()); @@ -308,7 +308,7 @@ void TestRemoveCopyIfStencilSimple() Vector result(5); typename Vector::iterator end = - thrust::remove_copy_if(data.begin(), data.end(), stencil.begin(), result.begin(), thrust::identity()); + 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..d5b7eba14b0 100644 --- a/thrust/testing/tabulate.cu +++ b/thrust/testing/tabulate.cu @@ -16,7 +16,7 @@ void TestTabulateDispatchExplicit() thrust::device_vector vec(1); my_system sys(0); - thrust::tabulate(sys, vec.begin(), vec.end(), thrust::identity()); + 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 vec(1); - thrust::tabulate(thrust::retag(vec.begin()), thrust::retag(vec.end()), thrust::identity()); + thrust::tabulate(thrust::retag(vec.begin()), thrust::retag(vec.end()), ::cuda::std::identity{}); ASSERT_EQUAL(13, vec.front()); } @@ -46,7 +46,7 @@ void TestTabulateSimple() Vector v(5); - thrust::tabulate(v.begin(), v.end(), thrust::identity()); + thrust::tabulate(v.begin(), v.end(), ::cuda::std::identity{}); Vector ref{0, 1, 2, 3, 4}; ASSERT_EQUAL(v, ref); @@ -88,7 +88,7 @@ void TestTabulateToDiscardIterator(size_t n) { thrust::tabulate(thrust::discard_iterator(), thrust::discard_iterator(n), - thrust::identity()); + ::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(), thrust::identity()); + iter = thrust::transform_if(input.begin(), input.end(), output.begin(), thrust::negate(), ::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(), thrust::identity()); + input.begin(), input.end(), stencil.begin(), output.begin(), thrust::negate(), ::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 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 h_result(n); thrust::device_vector d_result(n); - thrust::transform(h_first, h_first + n, h_result.begin(), thrust::identity()); - thrust::transform(d_first, d_first + n, d_result.begin(), thrust::identity()); + 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(), thrust::identity()), - thrust::make_transform_input_output_iterator(h_data.end(), thrust::negate(), thrust::identity()), + thrust::make_transform_input_output_iterator(h_data.begin(), thrust::negate(), ::cuda::std::identity{}), + thrust::make_transform_input_output_iterator(h_data.end(), thrust::negate(), ::cuda::std::identity{}), h_result.begin()); // run on device (uses reverse iterator negate) thrust::inclusive_scan( 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{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{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{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..d894c406382 100644 --- a/thrust/testing/type_traits.cu +++ b/thrust/testing/type_traits.cu @@ -38,7 +38,7 @@ void TestIsContiguousIterator() using ConstantIterator = thrust::constant_iterator; using CountingIterator = thrust::counting_iterator; - using TransformIterator = thrust::transform_iterator, HostVector::iterator>; + using TransformIterator = thrust::transform_iterator, HostVector::iterator>; // TODO using ZipIterator = thrust::zip_iterator; ASSERT_EQUAL((bool) thrust::is_contiguous_iterator::value, false); 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()); // 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 * 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()); + * 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()); + * thrust::transform(exec, data, data, + 4, data, ::cuda::std::identity{}); * * return 0; * } 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()); + * ::cuda::std::identity{}); * // indices now contains [1,2,5,7] * * return 0; diff --git a/thrust/thrust/logical.h b/thrust/thrust/logical.h index e10b352f7d5..409f23e6a8f 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()); // returns true - * thrust::all_of(thrust::host, A, A + 3, thrust::identity()); // 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()); // returns true + * thrust::all_of(thrust::host, A, A, ::cuda::std::identity{}); // returns true * * \endcode * @@ -104,11 +104,11 @@ all_of(const thrust::detail::execution_policy_base& exec, * ... * bool A[3] = {true, true, false}; * - * thrust::all_of(A, A + 2, thrust::identity()); // returns true - * thrust::all_of(A, A + 3, thrust::identity()); // 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()); // returns true + * thrust::all_of(A, A, ::cuda::std::identity{}); // returns true * * \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()); // returns true - * thrust::any_of(thrust::host, A, A + 3, thrust::identity()); // 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()); // 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()); // 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& exec, * ... * bool A[3] = {true, true, false}; * - * thrust::any_of(A, A + 2, thrust::identity()); // returns true - * thrust::any_of(A, A + 3, thrust::identity()); // 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()); // returns false + * thrust::any_of(A + 2, A + 3, ::cuda::std::identity{}); // returns false * * // empty range - * thrust::any_of(A, A, thrust::identity()); // 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()); // returns false - * thrust::none_of(thrust::host, A, A + 3, thrust::identity()); // 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()); // 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()); // 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& exec, * ... * bool A[3] = {true, true, false}; * - * thrust::none_of(A, A + 2, thrust::identity()); // returns false - * thrust::none_of(A, A + 3, thrust::identity()); // 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()); // returns true + * thrust::none_of(A + 2, A + 3, ::cuda::std::identity{}); // returns true * * // empty range - * thrust::none_of(A, A, thrust::identity()); // 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 b34c09b24db..f051612787e 100644 --- a/thrust/thrust/partition.h +++ b/thrust/thrust/partition.h @@ -485,7 +485,7 @@ thrust::pair 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()); + * 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 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()); + * 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 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()); + * 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 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()); + * 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 6a047632652..9cb56375638 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 *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 *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()); + * 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()); + * 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 04330e071f1..8ea61dd54b0 100644 --- a/thrust/thrust/system/cuda/detail/async/copy.h +++ b/thrust/thrust/system/cuda/detail/async/copy.h @@ -128,9 +128,7 @@ auto async_copy_n(thrust::cuda::execution_policy& from_exec, decltype(is_device_to_device_copy(from_exec, to_exec))>::value, unique_eager_event>::type { - using T = typename iterator_traits::value_type; - - return async_transform_n(select_device_system(from_exec, to_exec), first, n, output, thrust::identity()); + return async_transform_n(select_device_system(from_exec, to_exec), first, n, output, ::cuda::std::identity{}); } template 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 OutputIt THRUST_RUNTIME_FUNCTION device_to_device( execution_policy& policy, InputIt first, InputIt last, OutputIt result, thrust::detail::false_type) { - using InputTy = typename thrust::iterator_traits::value_type; - return cuda_cub::transform(policy, first, last, result, thrust::identity()); + return cuda_cub::transform(policy, first, last, result, ::cuda::std::identity{}); } template 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 # include -# include +# include 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..8df550d9b89 100644 --- a/thrust/thrust/system/detail/generic/copy.inl +++ b/thrust/thrust/system/detail/generic/copy.inl @@ -46,8 +46,7 @@ template & exec, InputIterator first, InputIterator last, OutputIterator result) { - using T = typename thrust::iterator_value::type; - return thrust::transform(exec, first, last, result, thrust::identity()); + return thrust::transform(exec, first, last, result, ::cuda::std::identity{}); } // end copy() template @@ -55,7 +54,7 @@ _CCCL_HOST_DEVICE OutputIterator copy_n(thrust::execution_policy& exec, InputIterator first, Size n, OutputIterator result) { using value_type = typename thrust::iterator_value::type; - using xfrm_type = thrust::identity; + using xfrm_type = thrust::identity; // TODO using functor_type = thrust::detail::unary_transform_functor; diff --git a/thrust/thrust/system/detail/generic/copy_if.inl b/thrust/thrust/system/detail/generic/copy_if.inl index 00075933e7c..660913fedc2 100644 --- a/thrust/thrust/system/detail/generic/copy_if.inl +++ b/thrust/thrust/system/detail/generic/copy_if.inl @@ -84,8 +84,7 @@ _CCCL_HOST_DEVICE OutputIterator copy_if( thrust::plus()); // scatter the true elements - thrust::scatter_if( - exec, first, last, scatter_indices.begin(), predicates.begin(), result, thrust::identity()); + 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::type>()); + ::cuda::std::identity{}); } // end gather() template ::type; - return thrust::gather_if(exec, map_first, map_last, stencil, input_first, result, thrust::identity()); + return thrust::gather_if(exec, map_first, map_last, stencil, input_first, result, ::cuda::std::identity{}); } // end gather_if() template ::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(), + ::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 #include #include #include #include -#include THRUST_NAMESPACE_BEGIN namespace system @@ -47,12 +47,7 @@ scatter(thrust::execution_policy& exec, InputIterator2 map, RandomAccessIterator output) { - thrust::transform( - exec, - first, - last, - thrust::make_permutation_iterator(output, map), - thrust::identity::type>()); + thrust::copy(exec, first, last, thrust::make_permutation_iterator(output, map)); } // end scatter() template ::type; - thrust::scatter_if(exec, first, last, map, stencil, output, thrust::identity()); + thrust::scatter_if(exec, first, last, map, stencil, output, ::cuda::std::identity{}); } // end scatter_if() template ::type; thrust::transform_if( - exec, first, last, stencil, thrust::make_permutation_iterator(output, map), thrust::identity(), 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 c1689eecddd..7353aef4e8a 100644 --- a/thrust/thrust/transform.h +++ b/thrust/thrust/transform.h @@ -326,7 +326,6 @@ transform(InputIterator1 first1, InputIterator1 last1, InputIterator2 first2, Ou * }; * * thrust::negate op; - * thrust::identity identity; * * // negate odd elements * thrust::transform_if(thrust::host, data, data + 10, data, op, is_odd()); // in-place transformation @@ -400,7 +399,6 @@ _CCCL_HOST_DEVICE ForwardIterator transform_if( * }; * * thrust::negate op; - * thrust::identity identity; * * // negate odd elements * thrust::transform_if(data, data + 10, data, op, is_odd()); // in-place transformation @@ -469,7 +467,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 op; - * thrust::identity identity; + * ::cuda::std::identity identity; * * thrust::transform_if(thrust::host, data, data + 10, stencil, data, op, identity); // in-place transformation * @@ -541,7 +539,7 @@ _CCCL_HOST_DEVICE ForwardIterator transform_if( * int stencil[10] = { 1, 0, 1, 0, 1, 0, 1, 0, 1, 0}; * * thrust::negate op; - * thrust::identity identity; + * ::cuda::std::identity identity; * * thrust::transform_if(data, data + 10, stencil, data, op, identity); // in-place transformation * @@ -623,7 +621,7 @@ ForwardIterator transform_if( * int output[6]; * * thrust::plus op; - * thrust::identity identity; + * ::cuda::std::identity identity; * * thrust::transform_if(thrust::host, input1, input1 + 6, input2, stencil, output, op, identity); * @@ -703,7 +701,7 @@ _CCCL_HOST_DEVICE ForwardIterator transform_if( * int output[6]; * * thrust::plus op; - * thrust::identity identity; + * ::cuda::std::identity identity; * * thrust::transform_if(input1, input1 + 6, input2, stencil, output, op, identity); *