Skip to content

Commit

Permalink
Backport to 2.8: Fix transform_iterator<identity> (#3652) and Depreca…
Browse files Browse the repository at this point in the history 
…te thrust::identity (#3649) (#3693)

* Fix transform_iterator<identity> and drop result_of_adaptable_function (#3652)

Thrust's transform_iterator relied on thrust::identity<T>::result_type via result_of_adaptable_function to avoid dangling references. However, thrust::identity<void>, cuda::std::identity, and the placeholder expression _1 have the same issue without a workaround. This change centralizes and adds workarounds for all of these by introducing a new trait thrust::detail::transform_iterator_reference, which decays the return value type of any of the aforementioned function objects.

* Deprecate and replace thrust::identity (#3649)
Co-authored-by: Michael Schellenberger Costa <[email protected]>
  • Loading branch information
@bernhardmgruber
bernhardmgruber authored Feb 5, 2025
1 parent 95919f9 commit 665f4f3
Show file tree
Hide file tree
Showing 44 changed files with 254 additions and 303 deletions.
2 changes: 1 addition & 1 deletion libcudacxx/examples/concurrent_hash_table.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
2 changes: 1 addition & 1 deletion thrust/examples/counting_iterator.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down
6 changes: 3 additions & 3 deletions thrust/examples/minimal_custom_backend.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
}
2 changes: 1 addition & 1 deletion thrust/testing/cuda/copy_if.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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());
Expand Down
12 changes: 6 additions & 6 deletions thrust/testing/cuda/is_partitioned.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand All @@ -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);
}
Expand Down
72 changes: 36 additions & 36 deletions thrust/testing/cuda/logical.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand All @@ -28,39 +28,39 @@ 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);
}

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);
}

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);
}

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);
}

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);
Expand Down Expand Up @@ -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);
}
Expand All @@ -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);
Expand All @@ -132,39 +132,39 @@ 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);
}

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);
}

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);
}

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);
}

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);
Expand Down Expand Up @@ -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);
}
Expand All @@ -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);
Expand All @@ -236,39 +236,39 @@ 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);
}

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);
}

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);
}

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);
}

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);
Expand Down Expand Up @@ -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);
}
Expand Down
4 changes: 2 additions & 2 deletions thrust/testing/cuda/partition_point.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
}
Expand Down
4 changes: 2 additions & 2 deletions thrust/testing/cuda/remove.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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());
Expand Down Expand Up @@ -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());
Expand Down
4 changes: 2 additions & 2 deletions thrust/testing/cuda/tabulate.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -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};
Expand Down
6 changes: 3 additions & 3 deletions thrust/testing/cuda/transform.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);

Expand Down Expand Up @@ -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);
Expand Down Expand Up @@ -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);

Expand Down
Loading

0 comments on commit 665f4f3

Please sign in to comment.