bgirgis/cpp-thirdparty
1
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

update boost on linux

Browse Source
This commit is contained in:
Bassem Girgis
2019-08-10 16:06:25 -05:00
parent 76ad52be58
commit 861b918727
5363 changed files with 483306 additions and 116507 deletions
Download Patch File Download Diff File Expand all files Collapse all files

45
linx64/include/boost/compute/algorithm/detail/copy_on_device.hpp
View File

@@ -32,7 +32,8 @@ template<class InputIterator, class OutputIterator>
inline event copy_on_device_cpu(InputIterator first,
OutputIterator result,
size_t count,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
meta_kernel k("copy");
const device& device = queue.get_device();
@@ -52,14 +53,15 @@ inline event copy_on_device_cpu(InputIterator first,
size_t global_work_size = device.compute_units();
if(count <= 1024) global_work_size = 1;
return k.exec_1d(queue, 0, global_work_size);
return k.exec_1d(queue, 0, global_work_size, events);
}
template<class InputIterator, class OutputIterator>
inline event copy_on_device_gpu(InputIterator first,
OutputIterator result,
size_t count,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename std::iterator_traits<InputIterator>::value_type input_type;
@@ -86,14 +88,15 @@ inline event copy_on_device_gpu(InputIterator first,
k.add_set_arg<const uint_>("count", static_cast<uint_>(count));
size_t global_work_size = calculate_work_size(count, vpt, tpb);
return k.exec_1d(queue, 0, global_work_size, tpb);
return k.exec_1d(queue, 0, global_work_size, tpb, events);
}
template<class InputIterator, class OutputIterator>
inline event dispatch_copy_on_device(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
const size_t count = detail::iterator_range_size(first, last);
@@ -108,18 +111,19 @@ inline event dispatch_copy_on_device(InputIterator first,
// See https://github.com/boostorg/compute/pull/626
if((device.type() & device::cpu) && !is_apple_platform_device(device))
{
return copy_on_device_cpu(first, result, count, queue);
return copy_on_device_cpu(first, result, count, queue, events);
}
return copy_on_device_gpu(first, result, count, queue);
return copy_on_device_gpu(first, result, count, queue, events);
}
template<class InputIterator, class OutputIterator>
inline OutputIterator copy_on_device(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
dispatch_copy_on_device(first, last, result, queue);
dispatch_copy_on_device(first, last, result, queue, events);
return result + std::distance(first, last);
}
@@ -127,9 +131,11 @@ template<class InputIterator>
inline discard_iterator copy_on_device(InputIterator first,
InputIterator last,
discard_iterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
(void) queue;
(void) events;
return result + std::distance(first, last);
}
@@ -138,19 +144,21 @@ template<class InputIterator, class OutputIterator>
inline future<OutputIterator> copy_on_device_async(InputIterator first,
InputIterator last,
OutputIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
event event_ = dispatch_copy_on_device(first, last, result, queue);
event event_ = dispatch_copy_on_device(first, last, result, queue, events);
return make_future(result + std::distance(first, last), event_);
}
#ifdef CL_VERSION_2_0
#ifdef BOOST_COMPUTE_CL_VERSION_2_0
// copy_on_device() specialization for svm_ptr
template<class T>
inline svm_ptr<T> copy_on_device(svm_ptr<T> first,
svm_ptr<T> last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -158,7 +166,7 @@ inline svm_ptr<T> copy_on_device(svm_ptr<T> first,
}
queue.enqueue_svm_memcpy(
result.get(), first.get(), count * sizeof(T)
result.get(), first.get(), count * sizeof(T), events
);
return result + count;
@@ -168,7 +176,8 @@ template<class T>
inline future<svm_ptr<T> > copy_on_device_async(svm_ptr<T> first,
svm_ptr<T> last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -176,12 +185,12 @@ inline future<svm_ptr<T> > copy_on_device_async(svm_ptr<T> first,
}
event event_ = queue.enqueue_svm_memcpy_async(
result.get(), first.get(), count * sizeof(T)
result.get(), first.get(), count * sizeof(T), events
);
return make_future(result + count, event_);
}
#endif // CL_VERSION_2_0
#endif // BOOST_COMPUTE_CL_VERSION_2_0
} // end detail namespace
} // end compute namespace

39
linx64/include/boost/compute/algorithm/detail/copy_to_device.hpp
View File

@@ -28,7 +28,8 @@ template<class HostIterator, class DeviceIterator>
inline DeviceIterator copy_to_device(HostIterator first,
HostIterator last,
DeviceIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -47,7 +48,8 @@ inline DeviceIterator copy_to_device(HostIterator first,
queue.enqueue_write_buffer(result.get_buffer(),
offset * sizeof(value_type),
count * sizeof(value_type),
::boost::addressof(*first));
::boost::addressof(*first),
events);
return result + static_cast<difference_type>(count);
}
@@ -56,7 +58,8 @@ template<class HostIterator, class DeviceIterator>
inline DeviceIterator copy_to_device_map(HostIterator first,
HostIterator last,
DeviceIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -78,7 +81,8 @@ inline DeviceIterator copy_to_device_map(HostIterator first,
result.get_buffer(),
CL_MAP_WRITE,
offset * sizeof(value_type),
count * sizeof(value_type)
count * sizeof(value_type),
events
)
);
@@ -99,7 +103,8 @@ template<class HostIterator, class DeviceIterator>
inline future<DeviceIterator> copy_to_device_async(HostIterator first,
HostIterator last,
DeviceIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -119,18 +124,20 @@ inline future<DeviceIterator> copy_to_device_async(HostIterator first,
queue.enqueue_write_buffer_async(result.get_buffer(),
offset * sizeof(value_type),
count * sizeof(value_type),
::boost::addressof(*first));
::boost::addressof(*first),
events);
return make_future(result + static_cast<difference_type>(count), event_);
}
#ifdef CL_VERSION_2_0
#ifdef BOOST_COMPUTE_CL_VERSION_2_0
// copy_to_device() specialization for svm_ptr
template<class HostIterator, class T>
inline svm_ptr<T> copy_to_device(HostIterator first,
HostIterator last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -138,7 +145,7 @@ inline svm_ptr<T> copy_to_device(HostIterator first,
}
queue.enqueue_svm_memcpy(
result.get(), ::boost::addressof(*first), count * sizeof(T)
result.get(), ::boost::addressof(*first), count * sizeof(T), events
);
return result + count;
@@ -148,7 +155,8 @@ template<class HostIterator, class T>
inline future<svm_ptr<T> > copy_to_device_async(HostIterator first,
HostIterator last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -156,7 +164,7 @@ inline future<svm_ptr<T> > copy_to_device_async(HostIterator first,
}
event event_ = queue.enqueue_svm_memcpy_async(
result.get(), ::boost::addressof(*first), count * sizeof(T)
result.get(), ::boost::addressof(*first), count * sizeof(T), events
);
return make_future(result + count, event_);
@@ -166,7 +174,8 @@ template<class HostIterator, class T>
inline svm_ptr<T> copy_to_device_map(HostIterator first,
HostIterator last,
svm_ptr<T> result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -174,7 +183,9 @@ inline svm_ptr<T> copy_to_device_map(HostIterator first,
}
// map
queue.enqueue_svm_map(result.get(), count * sizeof(T), CL_MAP_WRITE);
queue.enqueue_svm_map(
result.get(), count * sizeof(T), CL_MAP_WRITE, events
);
// copy [first; last) to result buffer
std::copy(first, last, static_cast<T*>(result.get()));
@@ -184,7 +195,7 @@ inline svm_ptr<T> copy_to_device_map(HostIterator first,
return result + count;
}
#endif // CL_VERSION_2_0
#endif // BOOST_COMPUTE_CL_VERSION_2_0
} // end detail namespace
} // end compute namespace

37
linx64/include/boost/compute/algorithm/detail/copy_to_host.hpp
View File

@@ -29,7 +29,8 @@ template<class DeviceIterator, class HostIterator>
inline HostIterator copy_to_host(DeviceIterator first,
DeviceIterator last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -46,7 +47,8 @@ inline HostIterator copy_to_host(DeviceIterator first,
queue.enqueue_read_buffer(buffer,
offset * sizeof(value_type),
count * sizeof(value_type),
::boost::addressof(*result));
::boost::addressof(*result),
events);
return iterator_plus_distance(result, count);
}
@@ -55,7 +57,8 @@ template<class DeviceIterator, class HostIterator>
inline HostIterator copy_to_host_map(DeviceIterator first,
DeviceIterator last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -77,7 +80,8 @@ inline HostIterator copy_to_host_map(DeviceIterator first,
first.get_buffer(),
CL_MAP_READ,
offset * sizeof(value_type),
count * sizeof(value_type)
count * sizeof(value_type),
events
)
);
@@ -102,7 +106,8 @@ template<class DeviceIterator, class HostIterator>
inline future<HostIterator> copy_to_host_async(DeviceIterator first,
DeviceIterator last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
typedef typename
std::iterator_traits<DeviceIterator>::value_type
@@ -120,18 +125,20 @@ inline future<HostIterator> copy_to_host_async(DeviceIterator first,
queue.enqueue_read_buffer_async(buffer,
offset * sizeof(value_type),
count * sizeof(value_type),
::boost::addressof(*result));
::boost::addressof(*result),
events);
return make_future(iterator_plus_distance(result, count), event_);
}
#ifdef CL_VERSION_2_0
#ifdef BOOST_COMPUTE_CL_VERSION_2_0
// copy_to_host() specialization for svm_ptr
template<class T, class HostIterator>
inline HostIterator copy_to_host(svm_ptr<T> first,
svm_ptr<T> last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -139,7 +146,7 @@ inline HostIterator copy_to_host(svm_ptr<T> first,
}
queue.enqueue_svm_memcpy(
::boost::addressof(*result), first.get(), count * sizeof(T)
::boost::addressof(*result), first.get(), count * sizeof(T), events
);
return result + count;
@@ -149,7 +156,8 @@ template<class T, class HostIterator>
inline future<HostIterator> copy_to_host_async(svm_ptr<T> first,
svm_ptr<T> last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -157,7 +165,7 @@ inline future<HostIterator> copy_to_host_async(svm_ptr<T> first,
}
event event_ = queue.enqueue_svm_memcpy_async(
::boost::addressof(*result), first.get(), count * sizeof(T)
::boost::addressof(*result), first.get(), count * sizeof(T), events
);
return make_future(iterator_plus_distance(result, count), event_);
@@ -167,7 +175,8 @@ template<class T, class HostIterator>
inline HostIterator copy_to_host_map(svm_ptr<T> first,
svm_ptr<T> last,
HostIterator result,
command_queue &queue)
command_queue &queue,
const wait_list &events)
{
size_t count = iterator_range_size(first, last);
if(count == 0){
@@ -175,7 +184,7 @@ inline HostIterator copy_to_host_map(svm_ptr<T> first,
}
// map
queue.enqueue_svm_map(first.get(), count * sizeof(T), CL_MAP_READ);
queue.enqueue_svm_map(first.get(), count * sizeof(T), CL_MAP_READ, events);
// copy [first; last) to result
std::copy(
@@ -189,7 +198,7 @@ inline HostIterator copy_to_host_map(svm_ptr<T> first,
return iterator_plus_distance(result, count);
}
#endif // CL_VERSION_2_0
#endif // BOOST_COMPUTE_CL_VERSION_2_0
} // end detail namespace
} // end compute namespace

2
linx64/include/boost/compute/algorithm/detail/find_extrema.hpp
View File

@@ -56,7 +56,7 @@ inline InputIterator find_extrema(InputIterator first,
// use serial method for OpenCL version 1.0 due to
// problems with atomic_cmpxchg()
#ifndef CL_VERSION_1_1
#ifndef BOOST_COMPUTE_CL_VERSION_1_1
return serial_find_extrema(first, last, compare, find_minimum, queue);
#endif

1
linx64/include/boost/compute/algorithm/detail/find_extrema_with_reduce.hpp
View File

@@ -246,6 +246,7 @@ inline void find_extrema_with_reduce(InputIterator input,
);
}
// Space complexity: \Omega(2 * work-group-size * work-groups-per-compute-unit)
template<class InputIterator, class Compare>
InputIterator find_extrema_with_reduce(InputIterator first,
InputIterator last,

1
linx64/include/boost/compute/algorithm/detail/find_if_with_atomics.hpp
View File

@@ -153,6 +153,7 @@ inline InputIterator find_if_with_atomics_multiple_vpt(InputIterator first,
return first + static_cast<difference_type>(index.read(queue));
}
// Space complexity: O(1)
template<class InputIterator, class UnaryPredicate>
inline InputIterator find_if_with_atomics(InputIterator first,
InputIterator last,

18
linx64/include/boost/compute/algorithm/detail/merge_sort_on_gpu.hpp
View File

@@ -91,6 +91,7 @@ inline size_t bitonic_block_sort(KeyIterator keys_first,
command_queue &queue)
{
typedef typename std::iterator_traits<KeyIterator>::value_type key_type;
typedef typename std::iterator_traits<ValueIterator>::value_type value_type;
meta_kernel k("bitonic_block_sort");
size_t count_arg = k.add_arg<const uint_>("count");
@@ -169,8 +170,12 @@ inline size_t bitonic_block_sort(KeyIterator keys_first,
k.decl<bool>("compare") << " = " <<
compare(k.var<key_type>("sibling_key"),
k.var<key_type>("my_key")) << ";\n" <<
k.decl<bool>("equal") << " = !(compare || " <<
compare(k.var<key_type>("my_key"),
k.var<key_type>("sibling_key")) << ");\n" <<
k.decl<bool>("swap") <<
" = compare ^ (sibling_idx < lid) ^ direction;\n" <<
"swap = equal ? false : swap;\n" <<
"my_key = swap ? sibling_key : my_key;\n";
if(sort_by_key)
{
@@ -219,8 +224,12 @@ inline size_t bitonic_block_sort(KeyIterator keys_first,
k.decl<bool>("compare") << " = " <<
compare(k.var<key_type>("sibling_key"),
k.var<key_type>("my_key")) << ";\n" <<
k.decl<bool>("equal") << " = !(compare || " <<
compare(k.var<key_type>("my_key"),
k.var<key_type>("sibling_key")) << ");\n" <<
k.decl<bool>("swap") <<
" = compare ^ (sibling_idx < lid);\n" <<
"swap = equal ? false : swap;\n" <<
"my_key = swap ? sibling_key : my_key;\n";
if(sort_by_key)
{
@@ -249,8 +258,11 @@ inline size_t bitonic_block_sort(KeyIterator keys_first,
k.var<key_type>("my_key") << ";\n";
if(sort_by_key)
{
k << values_first[k.var<const uint_>("gid")] << " = " <<
values_first[k.var<const uint_>("offset + my_index")] << ";\n";
k <<
k.decl<value_type>("my_value") << " = " <<
values_first[k.var<const uint_>("offset + my_index")] << ";\n" <<
"barrier(CLK_GLOBAL_MEM_FENCE);\n" <<
values_first[k.var<const uint_>("gid")] << " = my_value;\n";
}
k <<
// end if
@@ -418,7 +430,7 @@ inline void merge_blocks_on_gpu(KeyIterator keys_first,
");\n" <<
"left_idx = equal ? mid_idx + 1 : left_idx + 1;\n" <<
"right_idx = equal ? right_idx : mid_idx;\n" <<
"upper_key = equal ? upper_key : " <<
"upper_key = " <<
keys_first[k.var<const uint_>("left_idx")] << ";\n" <<
"}\n" <<
"}\n" <<

10
linx64/include/boost/compute/algorithm/detail/radix_sort.hpp
View File

@@ -17,6 +17,9 @@
#include <boost/type_traits/is_signed.hpp>
#include <boost/type_traits/is_floating_point.hpp>
#include <boost/mpl/and.hpp>
#include <boost/mpl/not.hpp>
#include <boost/compute/kernel.hpp>
#include <boost/compute/program.hpp>
#include <boost/compute/command_queue.hpp>
@@ -173,8 +176,10 @@ const char radix_sort_source[] =
" uint sum = 0;\n"
" for(uint i = 0; i < K2_BITS; i++){\n"
" uint x = global_offsets[i] + last_block_offsets[i];\n"
" mem_fence(CLK_GLOBAL_MEM_FENCE);\n" // work around the RX 500/Vega bug, see #811
" global_offsets[i] = sum;\n"
" sum += x;\n"
" mem_fence(CLK_GLOBAL_MEM_FENCE);\n" // work around the RX Vega bug, see #811
" }\n"
"}\n"
@@ -305,9 +310,12 @@ inline void radix_sort_impl(const buffer_iterator<T> first,
options << " -DASC";
}
// get type definition if it is a custom struct
std::string custom_type_def = boost::compute::type_definition<T2>() + "\n";
// load radix sort program
program radix_sort_program = cache->get_or_build(
cache_key, options.str(), radix_sort_source, context
cache_key, options.str(), custom_type_def + radix_sort_source, context
);
kernel count_kernel(radix_sort_program, "count");

10
linx64/include/boost/compute/algorithm/detail/scan_on_cpu.hpp
View File

@@ -76,10 +76,10 @@ inline OutputIterator scan_on_cpu(InputIterator first,
k.add_arg<output_type *>(memory_object::global_memory, "block_partial_sums");
k <<
"uint block = " <<
"(uint)ceil(((float)count)/(get_global_size(0) + 1));\n" <<
"uint block = (count + get_global_size(0))/(get_global_size(0) + 1);\n" <<
"uint index = get_global_id(0) * block;\n" <<
"uint end = min(count, index + block);\n";
"uint end = min(count, index + block);\n" <<
"if(index >= end) return;\n";
if(!exclusive){
k <<
@@ -154,11 +154,9 @@ inline OutputIterator scan_on_cpu(InputIterator first,
l.add_arg<output_type *>(memory_object::global_memory, "block_partial_sums");
l <<
"uint block = " <<
"(uint)ceil(((float)count)/(get_global_size(0) + 1));\n" <<
"uint block = (count + get_global_size(0))/(get_global_size(0) + 1);\n" <<
"uint index = block + get_global_id(0) * block;\n" <<
"uint end = min(count, index + block);\n" <<
k.decl<output_type>("sum") << " = block_partial_sums[0];\n" <<
"for(uint i = 0; i < get_global_id(0); i++) {\n" <<
"sum = " << op(k.var<output_type>("sum"),

2
linx64/include/boost/compute/algorithm/detail/scan_on_gpu.hpp
View File

@@ -108,7 +108,7 @@ public:
// store sum for the block
if(exclusive){
*this <<
"if(lid == block_size - 1){\n" <<
"if(lid == block_size - 1 && gid < count) {\n" <<
" block_sums[get_group_id(0)] = " <<
op(first[expr<cl_uint>("gid")], var<T>("scratch[lid]")) <<
";\n" <<

1
linx64/include/boost/compute/algorithm/detail/serial_reduce.hpp
View File

@@ -20,6 +20,7 @@ namespace boost {
namespace compute {
namespace detail {
// Space complexity: O(1)
template<class InputIterator, class OutputIterator, class BinaryFunction>
inline void serial_reduce(InputIterator first,
InputIterator last,

6
linx64/include/boost/compute/algorithm/detail/serial_reduce_by_key.hpp
View File

@@ -55,11 +55,9 @@ inline size_t serial_reduce_by_key(InputKeyIterator keys_first,
size_t result_size_arg = k.add_arg<uint_ *>(memory_object::global_memory,
"result_size");
convert<result_type> to_result_type;
k <<
k.decl<result_type>("result") <<
" = " << to_result_type(values_first[0]) << ";\n" <<
" = " << values_first[0] << ";\n" <<
k.decl<key_type>("previous_key") << " = " << keys_first[0] << ";\n" <<
k.decl<result_type>("value") << ";\n" <<
k.decl<key_type>("key") << ";\n" <<
@@ -70,7 +68,7 @@ inline size_t serial_reduce_by_key(InputKeyIterator keys_first,
values_result[0] << " = result;\n" <<
"for(ulong i = 1; i < count; i++) {\n" <<
" value = " << to_result_type(values_first[k.var<uint_>("i")]) << ";\n" <<
" value = " << values_first[k.var<uint_>("i")] << ";\n" <<
" key = " << keys_first[k.var<uint_>("i")] << ";\n" <<
" if (" << predicate(k.var<key_type>("previous_key"),
k.var<key_type>("key")) << ") {\n" <<