Thread Sanitizer failures during thrust::reduce

[siboehm]

This code:

#define THRUST_DEVICE_SYSTEM THRUST_DEVICE_SYSTEM_TBB

#include <iostream>
#include <thrust/device_vector.h>
#include <thrust/reduce.h>

int main()
{
    thrust::device_vector<float> vec(30, 1.0);

    float res = thrust::reduce(vec.begin(), vec.end(), 0.0f);

    std::cout << res << std::endl;
}

Results in tsan failures (I checked all 4 combinations of clang and gcc, OpenMP and TBB):

SUMMARY: ThreadSanitizer: data race (/home/simon/.cache/bazel/_bazel_simon/bc7bfb2fb5f4ce0fca17a46af902a577/external/gcc_toolchain_files/x86_64-buildroot-linux-gnu/sysroot/../lib64/libtsan.so.2+0x6a443) in operator delete(void*, unsigned long, std::align_val_t)

Full stracktrace

``` WARNING: ThreadSanitizer: data race (pid=1388483) Read of size 4 at 0x7ffe0b2eb960 by thread T1: #0 thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >::body(thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >&, tbb::detail::d0::split) ??:? (thrust_issue+0x1f236b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#1 tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>::execute(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1f2062) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#2 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::outermost_worker_waiter&) ??:? (thrust_issue+0x27201f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#3 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::outermost_worker_waiter&) ??:? (thrust_issue+0x264776) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#4 tbb::detail::r1::arena::process(tbb::detail::r1::thread_data&) ??:? (thrust_issue+0x26082b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#5 tbb::detail::r1::market::process(rml::job&) ??:? (thrust_issue+0x2a4530) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#6 tbb::detail::r1::rml::private_worker::run() ??:? (thrust_issue+0x2ab925) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#7 tbb::detail::r1::rml::private_worker::thread_routine(void*) ??:? (thrust_issue+0x2ab81a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) Previous write of size 4 at 0x7ffe0b2eb960 by main thread: #0 void thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >::operator()(tbb::detail::d1::blocked_range const&) ??:? (thrust_issue+0x1f3f92) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#1 tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>::run_body(tbb::detail::d1::blocked_range&) ??:? (thrust_issue+0x1f343a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#2 void tbb::detail::d1::dynamic_grainsize_mode >::work_balance, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>, tbb::detail::d1::blocked_range >(tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d1::blocked_range&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1f282b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#3 void tbb::detail::d1::partition_type_base::execute, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>, tbb::detail::d1::blocked_range >(tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d1::blocked_range&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1f24f5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#4 tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>::execute(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1f20af) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#5 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::external_waiter&) ??:? (thrust_issue+0x2b858f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#6 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::external_waiter&) ??:? (thrust_issue+0x2b5556) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#7 tbb::detail::r1::task_dispatcher::execute_and_wait(tbb::detail::d1::task*, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b45bf) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#8 tbb::detail::r1::execute_and_wait(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b4449) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#9 tbb::detail::d1::execute_and_wait(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1eb815) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#10 tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>::run(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >&, tbb::detail::d1::auto_partitioner const&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1f1b7e) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#11 tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>::run(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >&, tbb::detail::d1::auto_partitioner const&) ??:? (thrust_issue+0x1f19c6) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#12 void tbb::detail::d1::parallel_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus > >(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >&) ??:? (thrust_issue+0x1f16d6) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#13 float thrust::system::tbb::detail::reduce >, float, thrust::plus >(thrust::system::tbb::detail::execution_policy&, thrust::detail::normal_iterator >, thrust::detail::normal_iterator >, float, thrust::plus) ??:? (thrust_issue+0x1f14ea) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#14 float thrust::reduce >, float, thrust::plus >(thrust::detail::execution_policy_base const&, thrust::detail::normal_iterator >, thrust::detail::normal_iterator >, float, thrust::plus) ??:? (thrust_issue+0x1f1378) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#15 float thrust::system::detail::generic::reduce >, float>(thrust::execution_policy&, thrust::detail::normal_iterator >, thrust::detail::normal_iterator >, float) ??:? (thrust_issue+0x1f12ab) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#16 float thrust::reduce >, float>(thrust::detail::execution_policy_base const&, thrust::detail::normal_iterator >, thrust::detail::normal_iterator >, float) ??:? (thrust_issue+0x1f11a5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#17 float thrust::reduce >, float>(thrust::detail::normal_iterator >, thrust::detail::normal_iterator >, float) ??:? (thrust_issue+0x1e8161) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#18 main ??:? (thrust_issue+0x1e7f43) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) Location is stack of main thread. Location is global '' at 0x000000000000 ([stack]+0x1e960) Thread T1 (tid=1388880, running) created by main thread at: #0 pthread_create ??:? (thrust_issue+0x1581ab) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#1 tbb::detail::r1::rml::internal::thread_monitor::launch(void* (*)(void*), void*, unsigned long) ??:? (thrust_issue+0x2ad936) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#2 tbb::detail::r1::rml::private_worker::wake_or_launch() ??:? (thrust_issue+0x2ace3c) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#3 tbb::detail::r1::rml::private_server::wake_some(int) ??:? (thrust_issue+0x2ac283) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#4 tbb::detail::r1::rml::private_server::adjust_job_count_estimate(int) ??:? (thrust_issue+0x2ac322) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#5 tbb::detail::r1::market::adjust_demand(tbb::detail::r1::arena&, int, bool) ??:? (thrust_issue+0x2a43f0) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#6 void tbb::detail::r1::arena::advertise_new_work<(tbb::detail::r1::arena::new_work_type)0>() ??:? (thrust_issue+0x2b532b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#7 tbb::detail::r1::spawn_and_notify(tbb::detail::d1::task&, tbb::detail::r1::arena_slot*, tbb::detail::r1::arena*) task_dispatcher.cpp:? (thrust_issue+0x2b3d7f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#8 tbb::detail::r1::spawn(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b3cf8) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#9 tbb::detail::d1::spawn(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1edb83) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#10 tbb::detail::d1::auto_partition_type::spawn_task(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1edac1) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#11 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::spawn_self(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ed191) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#12 void tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::offer_work_impl, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d0::split&>(tbb::detail::d1::execution_data&, tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d0::split&) ??:? (thrust_issue+0x1ecf2f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#13 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::offer_work(tbb::detail::d0::split&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ec968) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#14 void tbb::detail::d1::partition_type_base::execute, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>, tbb::detail::d1::blocked_range >(tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d1::blocked_range&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ec2ee) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#15 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::execute(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ebce3) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#16 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::external_waiter&) ??:? (thrust_issue+0x2b858f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#17 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::external_waiter&) ??:? (thrust_issue+0x2b5556) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#18 tbb::detail::r1::task_dispatcher::execute_and_wait(tbb::detail::d1::task*, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b45bf) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#19 tbb::detail::r1::execute_and_wait(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b4449) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#20 tbb::detail::d1::execute_and_wait(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1eb815) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#21 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::run(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&, tbb::detail::d1::auto_partitioner const&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1eadde) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#22 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::run(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&, tbb::detail::d1::auto_partitioner const&) ??:? (thrust_issue+0x1eab96) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#23 void tbb::detail::d1::parallel_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > >(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&) ??:? (thrust_issue+0x1ea886) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#24 thrust::device_ptr thrust::system::tbb::detail::for_each_n, unsigned long, thrust::detail::host_generate_functor > >(thrust::system::tbb::detail::execution_policy&, thrust::device_ptr, unsigned long, thrust::detail::host_generate_functor >) ??:? (thrust_issue+0x1ea6fe) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#25 thrust::device_ptr thrust::for_each_n, unsigned long, thrust::detail::host_generate_functor > >(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, thrust::detail::host_generate_functor >) ??:? (thrust_issue+0x1ea541) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#26 thrust::device_ptr thrust::system::detail::generic::generate_n, unsigned long, thrust::detail::fill_functor >(thrust::execution_policy&, thrust::device_ptr, unsigned long, thrust::detail::fill_functor) ??:? (thrust_issue+0x1ea3a5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#27 thrust::device_ptr thrust::generate_n, unsigned long, thrust::detail::fill_functor >(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, thrust::detail::fill_functor) ??:? (thrust_issue+0x1ea220) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#28 thrust::device_ptr thrust::system::detail::generic::fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1ea121) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/cccl#676 thrust::device_ptr thrust::fill_n, unsigned long, float>(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1ea04c) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#30 thrust::device_ptr thrust::system::detail::generic::detail::uninitialized_fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&, thrust::detail::integral_constant) ??:? (thrust_issue+0x1e9f92) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#31 thrust::device_ptr thrust::system::detail::generic::uninitialized_fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9ec2) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/cccl#677 thrust::device_ptr thrust::uninitialized_fill_n, unsigned long, float>(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9dcc) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#33 _ZN6thrust6detail23allocator_traits_detail20fill_construct_rangeINS_16device_allocatorIfEENS_10device_ptrIfEEmfEENS0_10disable_ifIXsr31has_effectful_member_construct2IT_NS0_15pointer_elementIT0_E4typeET2_EE5valueEvE4typeERS8_SA_T1_RKSD_ ??:? (thrust_issue+0x1e9d15) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#34 void thrust::detail::fill_construct_range, thrust::device_ptr, unsigned long, float>(thrust::device_allocator&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9c5f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#35 thrust::detail::contiguous_storage >::uninitialized_fill_n(thrust::detail::normal_iterator >, unsigned long, float const&) ??:? (thrust_issue+0x1e98a7) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#36 thrust::detail::vector_base >::fill_init(unsigned long, float const&) ??:? (thrust_issue+0x1e86e4) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#37 thrust::detail::vector_base >::vector_base(unsigned long, float const&) ??:? (thrust_issue+0x1e842f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#38 thrust::device_vector >::device_vector(unsigned long, float const&) ??:? (thrust_issue+0x1e80b5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#39 main ??:? (thrust_issue+0x1e7ef9) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) SUMMARY: ThreadSanitizer: data race ??:? in thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >::body(thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >&, tbb::detail::d0::split) ================== ================== WARNING: ThreadSanitizer: data race (pid=1388483) Read of size 4 at 0x7b4000000a28 by thread T5: #0 thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >::body(thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >&, tbb::detail::d0::split) ??:? (thrust_issue+0x1f236b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#1 tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>::execute(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1f2062) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#2 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::outermost_worker_waiter&) ??:? (thrust_issue+0x27201f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#3 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::outermost_worker_waiter&) ??:? (thrust_issue+0x264776) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#4 tbb::detail::r1::arena::process(tbb::detail::r1::thread_data&) ??:? (thrust_issue+0x26082b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#5 tbb::detail::r1::market::process(rml::job&) ??:? (thrust_issue+0x2a4530) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#6 tbb::detail::r1::rml::private_worker::run() ??:? (thrust_issue+0x2ab925) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#7 tbb::detail::r1::rml::private_worker::thread_routine(void*) ??:? (thrust_issue+0x2ab81a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) Previous write of size 4 at 0x7b4000000a28 by thread T1: #0 void thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >::operator()(tbb::detail::d1::blocked_range const&) ??:? (thrust_issue+0x1f3f4a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#1 tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>::run_body(tbb::detail::d1::blocked_range&) ??:? (thrust_issue+0x1f343a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#2 void tbb::detail::d1::dynamic_grainsize_mode >::work_balance, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>, tbb::detail::d1::blocked_range >(tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d1::blocked_range&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1f282b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#3 void tbb::detail::d1::partition_type_base::execute, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>, tbb::detail::d1::blocked_range >(tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d1::blocked_range&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1f24f5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#4 tbb::detail::d1::start_reduce, thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >, tbb::detail::d1::auto_partitioner const>::execute(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1f20af) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#5 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::outermost_worker_waiter&) ??:? (thrust_issue+0x27201f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#6 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::outermost_worker_waiter&) ??:? (thrust_issue+0x264776) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#7 tbb::detail::r1::arena::process(tbb::detail::r1::thread_data&) ??:? (thrust_issue+0x26082b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#8 tbb::detail::r1::market::process(rml::job&) ??:? (thrust_issue+0x2a4530) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#9 tbb::detail::r1::rml::private_worker::run() ??:? (thrust_issue+0x2ab925) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#10 tbb::detail::r1::rml::private_worker::thread_routine(void*) ??:? (thrust_issue+0x2ab81a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) Location is heap block of size 256 at 0x7b4000000a00 allocated by main thread: #0 memalign ??:? (thrust_issue+0x15751a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#1 tbb::detail::r1::std_cache_aligned_allocate(unsigned long, unsigned long) allocator.cpp:? (thrust_issue+0x27eca5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#2 tbb::detail::r1::cache_aligned_allocate(unsigned long) ??:? (thrust_issue+0x27edeb) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#3 tbb::detail::r1::small_object_pool_impl::allocate_impl(tbb::detail::d1::small_object_pool*&, unsigned long) ??:? (thrust_issue+0x2b05fd) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#4 tbb::detail::r1::allocate(tbb::detail::d1::small_object_pool*&, unsigned long, tbb::detail::d1::execution_data const&) ??:? (thrust_issue+0x2b04a4) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#5 tbb::detail::d1::tree_node* tbb::detail::d1::small_object_allocator::new_object(tbb::detail::d1::execution_data&, tbb::detail::d1::node*&, int&&, tbb::detail::d1::small_object_allocator&) ??:? (thrust_issue+0x1ed091) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#6 void tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::offer_work_impl, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d0::split&>(tbb::detail::d1::execution_data&, tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d0::split&) ??:? (thrust_issue+0x1ecedd) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#7 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::offer_work(tbb::detail::d0::split&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ec968) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#8 void tbb::detail::d1::partition_type_base::execute, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>, tbb::detail::d1::blocked_range >(tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d1::blocked_range&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ec2ee) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#9 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::execute(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ebce3) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#10 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::external_waiter&) ??:? (thrust_issue+0x2b858f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#11 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::external_waiter&) ??:? (thrust_issue+0x2b5556) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#12 tbb::detail::r1::task_dispatcher::execute_and_wait(tbb::detail::d1::task*, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b45bf) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#13 tbb::detail::r1::execute_and_wait(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b4449) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#14 tbb::detail::d1::execute_and_wait(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1eb815) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#15 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::run(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&, tbb::detail::d1::auto_partitioner const&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1eadde) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#16 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::run(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&, tbb::detail::d1::auto_partitioner const&) ??:? (thrust_issue+0x1eab96) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#17 void tbb::detail::d1::parallel_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > >(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&) ??:? (thrust_issue+0x1ea886) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#18 thrust::device_ptr thrust::system::tbb::detail::for_each_n, unsigned long, thrust::detail::host_generate_functor > >(thrust::system::tbb::detail::execution_policy&, thrust::device_ptr, unsigned long, thrust::detail::host_generate_functor >) ??:? (thrust_issue+0x1ea6fe) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#19 thrust::device_ptr thrust::for_each_n, unsigned long, thrust::detail::host_generate_functor > >(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, thrust::detail::host_generate_functor >) ??:? (thrust_issue+0x1ea541) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#20 thrust::device_ptr thrust::system::detail::generic::generate_n, unsigned long, thrust::detail::fill_functor >(thrust::execution_policy&, thrust::device_ptr, unsigned long, thrust::detail::fill_functor) ??:? (thrust_issue+0x1ea3a5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#21 thrust::device_ptr thrust::generate_n, unsigned long, thrust::detail::fill_functor >(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, thrust::detail::fill_functor) ??:? (thrust_issue+0x1ea220) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#22 thrust::device_ptr thrust::system::detail::generic::fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1ea121) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#23 thrust::device_ptr thrust::fill_n, unsigned long, float>(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1ea04c) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#24 thrust::device_ptr thrust::system::detail::generic::detail::uninitialized_fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&, thrust::detail::integral_constant) ??:? (thrust_issue+0x1e9f92) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#25 thrust::device_ptr thrust::system::detail::generic::uninitialized_fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9ec2) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#26 thrust::device_ptr thrust::uninitialized_fill_n, unsigned long, float>(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9dcc) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#27 _ZN6thrust6detail23allocator_traits_detail20fill_construct_rangeINS_16device_allocatorIfEENS_10device_ptrIfEEmfEENS0_10disable_ifIXsr31has_effectful_member_construct2IT_NS0_15pointer_elementIT0_E4typeET2_EE5valueEvE4typeERS8_SA_T1_RKSD_ ??:? (thrust_issue+0x1e9d15) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#28 void thrust::detail::fill_construct_range, thrust::device_ptr, unsigned long, float>(thrust::device_allocator&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9c5f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/cccl#676 thrust::detail::contiguous_storage >::uninitialized_fill_n(thrust::detail::normal_iterator >, unsigned long, float const&) ??:? (thrust_issue+0x1e98a7) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#30 thrust::detail::vector_base >::fill_init(unsigned long, float const&) ??:? (thrust_issue+0x1e86e4) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#31 thrust::detail::vector_base >::vector_base(unsigned long, float const&) ??:? (thrust_issue+0x1e842f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/cccl#677 thrust::device_vector >::device_vector(unsigned long, float const&) ??:? (thrust_issue+0x1e80b5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#33 main ??:? (thrust_issue+0x1e7ef9) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) Thread T5 (tid=1388884, running) created by thread T1 at: #0 pthread_create ??:? (thrust_issue+0x1581ab) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#1 tbb::detail::r1::rml::internal::thread_monitor::launch(void* (*)(void*), void*, unsigned long) ??:? (thrust_issue+0x2ad936) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#2 tbb::detail::r1::rml::private_worker::wake_or_launch() ??:? (thrust_issue+0x2ace3c) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#3 tbb::detail::r1::rml::private_server::wake_some(int) ??:? (thrust_issue+0x2ac283) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#4 tbb::detail::r1::rml::private_server::propagate_chain_reaction() ??:? (thrust_issue+0x2ac714) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#5 tbb::detail::r1::rml::private_worker::run() ??:? (thrust_issue+0x2ab86a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#6 tbb::detail::r1::rml::private_worker::thread_routine(void*) ??:? (thrust_issue+0x2ab81a) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) Thread T1 (tid=1388880, running) created by main thread at: #0 pthread_create ??:? (thrust_issue+0x1581ab) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#1 tbb::detail::r1::rml::internal::thread_monitor::launch(void* (*)(void*), void*, unsigned long) ??:? (thrust_issue+0x2ad936) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#2 tbb::detail::r1::rml::private_worker::wake_or_launch() ??:? (thrust_issue+0x2ace3c) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#3 tbb::detail::r1::rml::private_server::wake_some(int) ??:? (thrust_issue+0x2ac283) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#4 tbb::detail::r1::rml::private_server::adjust_job_count_estimate(int) ??:? (thrust_issue+0x2ac322) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#5 tbb::detail::r1::market::adjust_demand(tbb::detail::r1::arena&, int, bool) ??:? (thrust_issue+0x2a43f0) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#6 void tbb::detail::r1::arena::advertise_new_work<(tbb::detail::r1::arena::new_work_type)0>() ??:? (thrust_issue+0x2b532b) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#7 tbb::detail::r1::spawn_and_notify(tbb::detail::d1::task&, tbb::detail::r1::arena_slot*, tbb::detail::r1::arena*) task_dispatcher.cpp:? (thrust_issue+0x2b3d7f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#8 tbb::detail::r1::spawn(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b3cf8) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#9 tbb::detail::d1::spawn(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1edb83) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#10 tbb::detail::d1::auto_partition_type::spawn_task(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1edac1) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#11 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::spawn_self(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ed191) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#12 void tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::offer_work_impl, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d0::split&>(tbb::detail::d1::execution_data&, tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d0::split&) ??:? (thrust_issue+0x1ecf2f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#13 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::offer_work(tbb::detail::d0::split&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ec968) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#14 void tbb::detail::d1::partition_type_base::execute, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>, tbb::detail::d1::blocked_range >(tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>&, tbb::detail::d1::blocked_range&, tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ec2ee) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#15 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::execute(tbb::detail::d1::execution_data&) ??:? (thrust_issue+0x1ebce3) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#16 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::external_waiter&) ??:? (thrust_issue+0x2b858f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#17 tbb::detail::d1::task* tbb::detail::r1::task_dispatcher::local_wait_for_all(tbb::detail::d1::task*, tbb::detail::r1::external_waiter&) ??:? (thrust_issue+0x2b5556) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#18 tbb::detail::r1::task_dispatcher::execute_and_wait(tbb::detail::d1::task*, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b45bf) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#19 tbb::detail::r1::execute_and_wait(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x2b4449) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#20 tbb::detail::d1::execute_and_wait(tbb::detail::d1::task&, tbb::detail::d1::task_group_context&, tbb::detail::d1::wait_context&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1eb815) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#21 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::run(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&, tbb::detail::d1::auto_partitioner const&, tbb::detail::d1::task_group_context&) ??:? (thrust_issue+0x1eadde) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#22 tbb::detail::d1::start_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > >, tbb::detail::d1::auto_partitioner const>::run(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&, tbb::detail::d1::auto_partitioner const&) ??:? (thrust_issue+0x1eab96) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#23 void tbb::detail::d1::parallel_for, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > >(tbb::detail::d1::blocked_range const&, thrust::system::tbb::detail::for_each_detail::body, unsigned long, thrust::detail::host_generate_functor > > const&) ??:? (thrust_issue+0x1ea886) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#24 thrust::device_ptr thrust::system::tbb::detail::for_each_n, unsigned long, thrust::detail::host_generate_functor > >(thrust::system::tbb::detail::execution_policy&, thrust::device_ptr, unsigned long, thrust::detail::host_generate_functor >) ??:? (thrust_issue+0x1ea6fe) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#25 thrust::device_ptr thrust::for_each_n, unsigned long, thrust::detail::host_generate_functor > >(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, thrust::detail::host_generate_functor >) ??:? (thrust_issue+0x1ea541) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#26 thrust::device_ptr thrust::system::detail::generic::generate_n, unsigned long, thrust::detail::fill_functor >(thrust::execution_policy&, thrust::device_ptr, unsigned long, thrust::detail::fill_functor) ??:? (thrust_issue+0x1ea3a5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#27 thrust::device_ptr thrust::generate_n, unsigned long, thrust::detail::fill_functor >(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, thrust::detail::fill_functor) ??:? (thrust_issue+0x1ea220) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#28 thrust::device_ptr thrust::system::detail::generic::fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1ea121) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/cccl#676 thrust::device_ptr thrust::fill_n, unsigned long, float>(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1ea04c) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#30 thrust::device_ptr thrust::system::detail::generic::detail::uninitialized_fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&, thrust::detail::integral_constant) ??:? (thrust_issue+0x1e9f92) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#31 thrust::device_ptr thrust::system::detail::generic::uninitialized_fill_n, unsigned long, float>(thrust::execution_policy&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9ec2) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/cccl#677 thrust::device_ptr thrust::uninitialized_fill_n, unsigned long, float>(thrust::detail::execution_policy_base const&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9dcc) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#33 _ZN6thrust6detail23allocator_traits_detail20fill_construct_rangeINS_16device_allocatorIfEENS_10device_ptrIfEEmfEENS0_10disable_ifIXsr31has_effectful_member_construct2IT_NS0_15pointer_elementIT0_E4typeET2_EE5valueEvE4typeERS8_SA_T1_RKSD_ ??:? (thrust_issue+0x1e9d15) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#34 void thrust::detail::fill_construct_range, thrust::device_ptr, unsigned long, float>(thrust::device_allocator&, thrust::device_ptr, unsigned long, float const&) ??:? (thrust_issue+0x1e9c5f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#35 thrust::detail::contiguous_storage >::uninitialized_fill_n(thrust::detail::normal_iterator >, unsigned long, float const&) ??:? (thrust_issue+0x1e98a7) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#36 thrust::detail::vector_base >::fill_init(unsigned long, float const&) ??:? (thrust_issue+0x1e86e4) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#37 thrust::detail::vector_base >::vector_base(unsigned long, float const&) ??:? (thrust_issue+0x1e842f) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#38 thrust::device_vector >::device_vector(unsigned long, float const&) ??:? (thrust_issue+0x1e80b5) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) NVIDIA/thrust#39 main ??:? (thrust_issue+0x1e7ef9) (BuildId: 86ab2421a7f27167158b10d1a4c9d7c5) SUMMARY: ThreadSanitizer: data race ??:? in thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >::body(thrust::system::tbb::detail::reduce_detail::body >, float, thrust::plus >&, tbb::detail::d0::split) ================== 30 ThreadSanitizer: reported 2 warnings ```

The result is still correct, so I do not know if this is an actual bug (the llvm docs however state that there are not really any false positives). The practical impact is that this makes it impossible to use tsan in a project that utilizes thrust.

[gevtushenko]

@siboehm thank you for reporting this! I can reproduce the race. It's mostly caused by the split constructor:

struct body
{
  body(body& b, ::tbb::split)
    : first(b.first), sum(b.sum) /* <------ this read */, first_call(true), binary_op(b.binary_op)
  {}

  template <typename Size>
  void operator()(const ::tbb::blocked_range<Size> &r)
  {
     // ....
     sum = temp; // <------ races with this write
  } // end operator()()
}; // end body

According to the TBB docs:

The splitting constructor might run concurrently while object x is being used for the first half of the reduction.

In other words, it does seem like an actual race rather than a false positive. We also should accept constant references in join. We could rewrite the body in the canonical form, although it requires the accumulator type to be default constructible:

template <typename RandomAccessIterator, typename OutputType,
          typename BinaryFunction>
struct body {
  RandomAccessIterator first;
  OutputType sum;
  thrust::detail::wrapped_function<BinaryFunction, OutputType> binary_op;

  body(RandomAccessIterator first, BinaryFunction binary_op)
      : first(first)
      , sum{}
      , binary_op(binary_op) {}

  body(const body &b, ::tbb::split)
      : first(b.first)
      , sum{}
      , binary_op(b.binary_op) {}

  template <typename Size>
  void operator()(const ::tbb::blocked_range<Size> &r) {
    RandomAccessIterator iter = first + r.begin();

    OutputType temp = sum;

    for (Size i = r.begin(); i != r.end(); ++i, ++iter)
      temp = binary_op(temp, *iter);

    sum = temp;
  }

  void join(const body &b) { sum = binary_op(sum, b.sum); }
};