Add `thrust::get_per_device_resource` and `thrust::per_device_allocator`, which
the new asynchronous algorithms backend needs, because you don't want to use
a caching allocator that caches memory from another device.
Bug
ADD
thrust/per_device_resource.h
+/*
+ * Copyright 2008-2013 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/cpp11_required.h>
+
+#if THRUST_CPP_DIALECT >= 2011
+
+#include <thrust/system/detail/generic/per_device_resource.h>
+#include <thrust/system/detail/adl/per_device_resource.h>
+
+THRUST_BEGIN_NS
+
+/*! Returns a global instance of \p MR for the current device of the provided system.
+ *
+ * \tparam MR type of a memory resource to get an instance from. Must be \p DefaultConstructible.
+ * \param system execution policy for which the resource is requested.
+ * \returns a pointer to a global instance of \p MR for the current device.
+ */
+template<typename MR, typename DerivedPolicy>
+__host__
+MR * get_per_device_resource(const thrust::detail::execution_policy_base<DerivedPolicy> & system)
+{
+ using thrust::system::detail::generic::get_per_device_resource;
+
+ return get_per_device_resource<MR>(
+ thrust::detail::derived_cast(
+ thrust::detail::strip_const(system)));
+}
+
+/*! A helper allocator class that uses global per device instances of a given upstream memory resource. Requires the memory
+ * resource to be default constructible.
+ *
+ * \tparam T the type that will be allocated by this allocator.
+ * \tparam MR the upstream memory resource to use for memory allocation. Must derive from
+ * \p thrust::mr::memory_resource and must be \p final.
+ * \tparam ExecutionPolicy the execution policy of the system to be used to retrieve the resource for the current device.
+ */
+template<typename T, typename Upstream, typename ExecutionPolicy>
+class per_device_allocator : public thrust::mr::allocator<T, Upstream>
+{
+ typedef thrust::mr::allocator<T, Upstream> base;
+
+public:
+ /*! The \p rebind metafunction provides the type of an \p per_device_allocator instantiated with another type.
+ *
+ * \tparam U the other type to use for instantiation.
+ */
+ template<typename U>
+ struct rebind
+ {
+ /*! The typedef \p other gives the type of the rebound \p per_device_allocator.
+ */
+ typedef per_device_allocator<U, Upstream, ExecutionPolicy> other;
+ };
+
+ /*! Default constructor. Uses \p get_global_resource to get the global instance of \p Upstream and initializes the
+ * \p allocator base subobject with that resource.
+ */
+ __host__
+ per_device_allocator() : base(get_per_device_resource<Upstream>(ExecutionPolicy()))
+ {
+ }
+
+ /*! Copy constructor. Copies the memory resource pointer. */
+ __host__ __device__
+ per_device_allocator(const per_device_allocator & other)
+ : base(other) {}
+
+ /*! Conversion constructor from an allocator of a different type. Copies the memory resource pointer. */
+ template<typename U>
+ __host__ __device__
+ per_device_allocator(const per_device_allocator<U, Upstream, ExecutionPolicy> & other)
+ : base(other) {}
+
+ /*! Destructor. */
+ __host__ __device__
+ ~per_device_allocator() {}
+};
+
+
+THRUST_END_NS
+
+#endif // THRUST_CPP_DIALECT >= 2011

ADD
thrust/system/cpp/detail/per_device_resource.h
+/*
+ * Copyright 2018 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// this system has no special per device resource functions
+

MODIFY
thrust/system/cuda/detail/async/customization.h
thrust/system/cuda/detail/async/customization.h
#include <thrust/mr/allocator.h>
#include <thrust/mr/disjoint_sync_pool.h>
#include <thrust/mr/sync_pool.h>
+#include <thrust/per_device_resource.h>
THRUST_BEGIN_NS
@@ -80,8 +81,8 @@ auto get_async_device_allocator(
thrust::detail::execution_policy_base<DerivedPolicy>&
)
THRUST_DECLTYPE_RETURNS(
- thrust::mr::stateless_resource_allocator<
- thrust::detail::uint8_t, default_async_device_resource
+ thrust::per_device_allocator<
+ thrust::detail::uint8_t, default_async_device_resource, par_t
>{}
)

ADD
thrust/system/cuda/detail/per_device_resource.h
+/******************************************************************************
+ * Copyright (c) 2018, NVIDIA CORPORATION. All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ * * Redistributions of source code must retain the above copyright
+ * notice, this list of conditions and the following disclaimer.
+ * * Redistributions in binary form must reproduce the above copyright
+ * notice, this list of conditions and the following disclaimer in the
+ * documentation and/or other materials provided with the distribution.
+ * * Neither the name of the NVIDIA CORPORATION nor the
+ * names of its contributors may be used to endorse or promote products
+ * derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
+ * AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
+ * IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
+ * ARE DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/detail/cpp11_required.h>
+
+#if THRUST_CPP_DIALECT >= 2011
+
+#if THRUST_DEVICE_COMPILER == THRUST_DEVICE_COMPILER_NVCC
+
+#include <thrust/system/cuda/config.h>
+
+#include <thrust/system/cuda/detail/execution_policy.h>
+
+#include <mutex>
+#include <unordered_map>
+
+THRUST_BEGIN_NS
+
+namespace cuda_cub
+{
+
+template<typename MR, typename DerivedPolicy>
+__host__
+MR * get_per_device_resource(execution_policy<DerivedPolicy>&)
+{
+ static std::mutex map_lock;
+ static std::unordered_map<int, MR> device_id_to_resource;
+
+ int device_id;
+ thrust::cuda_cub::throw_on_error(cudaGetDevice(&device_id));
+
+ std::lock_guard<std::mutex> lock{map_lock};
+ return &device_id_to_resource[device_id];
+}
+
+}
+
+THRUST_END_NS
+
+#endif
+
+#endif
+

ADD
thrust/system/detail/adl/per_device_resource.h
+/*
+ * Copyright 2018 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// the purpose of this header is to #include the per_device_resource.h header
+// of the sequential, host, and device systems. It should be #included in any
+// code which uses adl to dispatch per_device_resource
+
+#include <thrust/system/detail/sequential/per_device_resource.h>
+
+#if 0
+#include <thrust/system/cpp/detail/per_device_resource.h>
+#include <thrust/system/cuda/detail/per_device_resource.h>
+#include <thrust/system/omp/detail/per_device_resource.h>
+#include <thrust/system/tbb/detail/per_device_resource.h>
+#endif
+
+#define __THRUST_HOST_SYSTEM_PER_DEVICE_RESOURCE_HEADER <__THRUST_HOST_SYSTEM_ROOT/detail/per_device_resource.h>
+#include __THRUST_HOST_SYSTEM_PER_DEVICE_RESOURCE_HEADER
+#undef __THRUST_HOST_SYSTEM_PER_DEVICE_RESOURCE_HEADER
+
+#define __THRUST_DEVICE_SYSTEM_PER_DEVICE_RESOURCE_HEADER <__THRUST_DEVICE_SYSTEM_ROOT/detail/per_device_resource.h>
+#include __THRUST_DEVICE_SYSTEM_PER_DEVICE_RESOURCE_HEADER
+#undef __THRUST_DEVICE_SYSTEM_PER_DEVICE_RESOURCE_HEADER
+

ADD
thrust/system/detail/generic/per_device_resource.h
+/*
+ * Copyright 2008-2013 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+#include <thrust/system/detail/generic/tag.h>
+#include <thrust/mr/memory_resource.h>
+
+namespace thrust
+{
+namespace system
+{
+namespace detail
+{
+namespace generic
+{
+
+
+template<typename MR, typename DerivedPolicy>
+__host__
+MR * get_per_device_resource(thrust::detail::execution_policy_base<DerivedPolicy>&)
+{
+ return mr::get_global_resource<MR>();
+}
+
+
+} // end generic
+} // end detail
+} // end system
+} // end thrust
+

ADD
thrust/system/detail/sequential/per_device_resource.h
+/*
+ * Copyright 2018 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// this system has no special per device resource functions
+

ADD
thrust/system/omp/detail/per_device_resource.h
+/*
+ * Copyright 2018 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// this system has no special per device resource functions
+

ADD
thrust/system/tbb/detail/per_device_resource.h
+/*
+ * Copyright 2018 NVIDIA Corporation
+ *
+ * Licensed under the Apache License, Version 2.0 (the "License");
+ * you may not use this file except in compliance with the License.
+ * You may obtain a copy of the License at
+ *
+ * http://www.apache.org/licenses/LICENSE-2.0
+ *
+ * Unless required by applicable law or agreed to in writing, software
+ * distributed under the License is distributed on an "AS IS" BASIS,
+ * WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+ * See the License for the specific language governing permissions and
+ * limitations under the License.
+ */
+
+#pragma once
+
+#include <thrust/detail/config.h>
+
+// this system has no special per device resource functions
+

Add `thrust::get_per_device_resource` and `thrust::per_device_allocator`, which
the new asynchronous algorithms backend needs, because you don't want to use
a caching allocator that caches memory from another device.

Bug 2455945