`thrust::async`:
* Add missing `.after` dependency extraction.
* Add the missing ref qualifier to exec_with_alloc to make it usable.
Bug
Bug
MODIFY
testing/allocator_aware_policies.cu
testing/allocator_aware_policies.cu
@@ -70,15 +70,45 @@ struct TestAllocatorAttachment
>::value), true);
}
+ template<typename Policy>
+ void test_temporary_allocation_valid(Policy policy)
+ {
+ using thrust::detail::get_temporary_buffer;
+
+ return_temporary_buffer(
+ policy,
+ get_temporary_buffer<int>(
+ policy,
+ 123
+ ).first
+ );
+ }
+
void operator()()
{
typename PolicyInfo::policy policy;
+ // test correctness of attachment
assert_correct<test_allocator_t<int> >(policy(test_allocator_t<int>()));
assert_correct<test_allocator_t<int>&>(policy(test_allocator));
assert_correct<test_allocator_t<int> >(policy(const_test_allocator));
assert_npa_correct<test_memory_resource_t>(policy(&test_memory_resource));
+
+ // test whether the resulting policy is actually usable
+ // a real allocator is necessary here, unlike above
+ std::allocator<int> alloc;
+ const std::allocator<int> const_alloc;
+
+ test_temporary_allocation_valid(policy(std::allocator<int>()));
+ test_temporary_allocation_valid(policy(alloc));
+ test_temporary_allocation_valid(policy(const_alloc));
+
+ #if THRUST_CPP_DIALECT >= 2011
+ test_temporary_allocation_valid(policy(std::allocator<int>()).after(1));
+ test_temporary_allocation_valid(policy(alloc).after(1));
+ test_temporary_allocation_valid(policy(const_alloc).after(1));
+ #endif
}
};

MODIFY
thrust/detail/execute_with_allocator.h
thrust/detail/execute_with_allocator.h
@@ -55,12 +55,13 @@ public:
: super_t(super), alloc(alloc_)
{}
+ __thrust_exec_check_disable__
__host__ __device__
execute_with_allocator(Allocator alloc_)
: alloc(alloc_)
{}
- Allocator get_allocator() { return alloc; }
+ typename remove_reference<Allocator>::type& get_allocator() { return alloc; }
#if __cplusplus >= 201103L
template<typename ...Dependencies>

MODIFY
thrust/detail/execute_with_dependencies.h
thrust/detail/execute_with_dependencies.h
@@ -104,7 +104,7 @@ public:
return std::move(dependencies);
}
- Allocator
+ typename remove_reference<Allocator>::type&
__host__
get_allocator()
{

MODIFY
thrust/system/cuda/detail/async/copy.h
thrust/system/cuda/detail/async/copy.h
@@ -392,10 +392,15 @@ auto async_copy_n(
fp = depend_on<return_type, return_pointer>(
[] (decltype(content) const& c)
{ return c.get(); }
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
std::move(content)
, unique_stream(nonowning, user_raw_stream)
)
+ , extract_dependencies(
+ std::move(select_device_system(from_exec, to_exec))
+ )
+ )
);
}
else
@@ -403,9 +408,14 @@ auto async_copy_n(
fp = depend_on<return_type, return_pointer>(
[] (decltype(content) const& c)
{ return c.get(); }
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
std::move(content)
)
+ , extract_dependencies(
+ std::move(select_device_system(from_exec, to_exec))
+ )
+ )
);
}

MODIFY
thrust/system/cuda/detail/async/for_each.h
thrust/system/cuda/detail/async/for_each.h
@@ -95,16 +95,23 @@ auto async_for_each_n(
{
fp = depend_on<void, pointer>(
nullptr
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
unique_stream(nonowning, user_raw_stream)
)
+ , extract_dependencies(
+ std::move(policy)
+ )
+ )
);
}
else
{
fp = depend_on<void, pointer>(
nullptr
- , std::make_tuple()
+ , extract_dependencies(
+ std::move(policy)
+ )
);
}

MODIFY
thrust/system/cuda/detail/async/reduce.h
thrust/system/cuda/detail/async/reduce.h
@@ -136,7 +136,7 @@ auto async_reduce_n(
, unique_stream(nonowning, user_raw_stream)
)
, extract_dependencies(
- std::move(thrust::detail::derived_cast(policy))
+ std::move(policy)
)
)
);
@@ -157,7 +157,7 @@ auto async_reduce_n(
std::move(content)
)
, extract_dependencies(
- std::move(thrust::detail::derived_cast(policy))
+ std::move(policy)
)
)
);

MODIFY
thrust/system/cuda/detail/async/sort.h
thrust/system/cuda/detail/async/sort.h
@@ -161,19 +161,29 @@ auto async_stable_sort_n(
{
fp = depend_on<void, pointer>(
nullptr
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
std::move(content)
, unique_stream(nonowning, user_raw_stream)
)
+ , extract_dependencies(
+ std::move(policy)
+ )
+ )
);
}
else
{
fp = depend_on<void, pointer>(
nullptr
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
std::move(content)
)
+ , extract_dependencies(
+ std::move(policy)
+ )
+ )
);
}
@@ -294,19 +304,29 @@ auto async_stable_sort_n(
{
fp = depend_on<void, pointer>(
nullptr
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
std::move(content)
, unique_stream(nonowning, user_raw_stream)
)
+ , extract_dependencies(
+ std::move(policy)
+ )
+ )
);
}
else
{
fp = depend_on<void, pointer>(
nullptr
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
std::move(content)
)
+ , extract_dependencies(
+ std::move(policy)
+ )
+ )
);
}

MODIFY
thrust/system/cuda/detail/async/transform.h
thrust/system/cuda/detail/async/transform.h
@@ -115,10 +115,15 @@ auto async_transform_n(
fp = depend_on<return_type, return_pointer>(
[] (decltype(content) const& c)
{ return c.get(); }
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
std::move(content)
, unique_stream(nonowning, user_raw_stream)
)
+ , extract_dependencies(
+ std::move(policy)
+ )
+ )
);
}
else
@@ -126,9 +131,14 @@ auto async_transform_n(
fp = depend_on<return_type, return_pointer>(
[] (decltype(content) const& c)
{ return c.get(); }
- , std::make_tuple(
+ , std::tuple_cat(
+ std::make_tuple(
std::move(content)
)
+ , extract_dependencies(
+ std::move(policy)
+ )
+ )
);
}

`thrust::async`:

* Add missing `.after` dependency extraction.
* Add the missing ref qualifier to exec_with_alloc to make it usable.

Bug 2379510
Bug 2379513