Skip to content

Develop thrust stream #61

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
neon60 merged 4 commits into from
May 6, 2019
Merged

Develop thrust stream #61

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
c8a7f13
Limit block size of device sorts so large types can fit in shared memory
ex-rzr Apr 30, 2019
e68200f
Add workaround for mem access fault in device scan on ROCm 2.3
ex-rzr May 1, 2019
2ae10b3
Merge branch 'support_very_large_types' into 'develop_thrust_stream'
neon60 May 1, 2019
7168e8d
Merge branch 'workaround_scan' into 'develop_thrust_stream'
neon60 May 2, 2019
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
3 changes: 2 additions & 1 deletion rocprim/include/rocprim/block/detail/block_sort_bitonic.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -235,11 +235,12 @@ class block_sort_bitonic
return !r;
};
wsort.sort(kv..., compare_function2);

#pragma unroll
for(unsigned int length = ::rocprim::warp_size(); length < Size; length *= 2)
{
bool dir = (flat_tid & (length * 2)) != 0;
#pragma unroll
for(unsigned int k = length; k > 0; k /= 2)
{
copy_to_shared(kv..., flat_tid, storage);
Expand Down
33 changes: 33 additions & 0 deletions rocprim/include/rocprim/device/config_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -52,6 +52,39 @@ struct kernel_config
namespace detail
{

template<
unsigned int MaxBlockSize,
unsigned int SharedMemoryPerThread,
// Most kernels require block sizes not smaller than warp
unsigned int MinBlockSize = ::rocprim::warp_size(),
// Can fit in shared memory?
// Although GPUs have 64KiB, 32KiB is used here as a "soft" limit,
// because some additional memory may be required in kernels
bool = (MaxBlockSize * SharedMemoryPerThread <= (1u << 15))
>
struct limit_block_size
{
// No, then try to decrease block size
static constexpr unsigned int value =
limit_block_size<
detail::next_power_of_two(MaxBlockSize) / 2,
SharedMemoryPerThread,
MinBlockSize
>::value;
};

template<
unsigned int MaxBlockSize,
unsigned int SharedMemoryPerThread,
unsigned int MinBlockSize
>
struct limit_block_size<MaxBlockSize, SharedMemoryPerThread, MinBlockSize, true>
{
static_assert(MaxBlockSize >= MinBlockSize, "Data is too large, it cannot fit in shared memory");

static constexpr unsigned int value = MaxBlockSize;
};

template<class...>
using void_t = void;

Expand Down
11 changes: 6 additions & 5 deletions rocprim/include/rocprim/device/detail/device_binary_search.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,6 +27,7 @@ namespace detail
{

template<class Size>
ROCPRIM_DEVICE inline
Size get_binary_search_middle(Size left, Size right)
{
// Instead of `/ 2` we use `* 33 / 64`, i.e. the middle is slightly moved.
Expand All @@ -40,7 +41,7 @@ Size get_binary_search_middle(Size left, Size right)
}

template<class RandomAccessIterator, class Size, class T, class BinaryPredicate>
ROCPRIM_DEVICE
ROCPRIM_DEVICE inline
Size lower_bound_n(RandomAccessIterator first,
Size size,
const T& value,
Expand All @@ -64,7 +65,7 @@ Size lower_bound_n(RandomAccessIterator first,
}

template<class RandomAccessIterator, class Size, class T, class BinaryPredicate>
ROCPRIM_DEVICE
ROCPRIM_DEVICE inline
Size upper_bound_n(RandomAccessIterator first,
Size size,
const T& value,
Expand All @@ -90,7 +91,7 @@ Size upper_bound_n(RandomAccessIterator first,
struct lower_bound_search_op
{
template<class HaystackIterator, class CompareOp, class Size, class T>
ROCPRIM_DEVICE
ROCPRIM_DEVICE inline
Size operator()(HaystackIterator haystack, Size size, const T& value, CompareOp compare_op) const
{
return lower_bound_n(haystack, size, value, compare_op);
Expand All @@ -100,7 +101,7 @@ struct lower_bound_search_op
struct upper_bound_search_op
{
template<class HaystackIterator, class CompareOp, class Size, class T>
ROCPRIM_DEVICE
ROCPRIM_DEVICE inline
Size operator()(HaystackIterator haystack, Size size, const T& value, CompareOp compare_op) const
{
return upper_bound_n(haystack, size, value, compare_op);
Expand All @@ -110,7 +111,7 @@ struct upper_bound_search_op
struct binary_search_op
{
template<class HaystackIterator, class CompareOp, class Size, class T>
ROCPRIM_DEVICE
ROCPRIM_DEVICE inline
bool operator()(HaystackIterator haystack, Size size, const T& value, CompareOp compare_op) const
{
const Size n = lower_bound_n(haystack, size, value, compare_op);
Expand Down
29 changes: 18 additions & 11 deletions rocprim/include/rocprim/device/detail/lookback_scan_state.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -323,21 +323,28 @@ class lookback_scan_prefix_op
T get_prefix()
{
flag_type flag;
T partial_prefix;
unsigned int previous_block_id = block_id_ - ::rocprim::lane_id() - 1;
bool is_prefix_initialized = false;
T prefix;

// reduce last warp_size() number of prefixes to
// get the complete prefix for this block.
reduce_partial_prefixes(previous_block_id, flag, partial_prefix);
T prefix = partial_prefix;

// while we don't load a complete prefix, reduce partial prefixes
while(::rocprim::detail::warp_all(flag != PREFIX_COMPLETE))
do
{
previous_block_id -= ::rocprim::warp_size();
// reduce last warp_size() number of prefixes to
// get the complete prefix for this block.
T partial_prefix;
reduce_partial_prefixes(previous_block_id, flag, partial_prefix);
prefix = scan_op_(partial_prefix, prefix);
}
if(!is_prefix_initialized)
{
prefix = partial_prefix;
is_prefix_initialized = true;
}
else
{
prefix = scan_op_(partial_prefix, prefix);
}
previous_block_id -= ::rocprim::warp_size();
// while we don't load a complete prefix, reduce partial prefixes
} while(::rocprim::detail::warp_all(flag != PREFIX_COMPLETE));
return prefix;
}

Expand Down
27 changes: 4 additions & 23 deletions rocprim/include/rocprim/device/device_merge_sort_config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -42,47 +42,28 @@ using merge_sort_config = kernel_config<BlockSize, 1>;
namespace detail
{

// TODO investigate why some tests fail with block size > 256
template<class Key, class Value>
struct merge_sort_config_803
{
// static constexpr size_t key_value_size = sizeof(Key) + sizeof(Value);
// static constexpr unsigned int item_scale =
// ::rocprim::detail::ceiling_div<unsigned int>(key_value_size, 8);

// using type = merge_sort_config<::rocprim::max(256U, 1024U / item_scale)>;
using type = merge_sort_config<256U>;
using type = merge_sort_config<limit_block_size<256U, sizeof(Key) + sizeof(Value)>::value>;
};

template<class Key>
struct merge_sort_config_803<Key, empty_type>
{
// static constexpr unsigned int item_scale =
// ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Key), 8);

// using type = merge_sort_config<::rocprim::max(256U, 1024U / item_scale)>;
using type = merge_sort_config<256U>;
using type = merge_sort_config<limit_block_size<256U, sizeof(Key)>::value>;
};

template<class Key, class Value>
struct merge_sort_config_900
{
// static constexpr size_t key_value_size = sizeof(Key) + sizeof(Value);
// static constexpr unsigned int item_scale =
// ::rocprim::detail::ceiling_div<unsigned int>(key_value_size, 16);

// using type = merge_sort_config<::rocprim::max(256U, 1024U / item_scale)>;
using type = merge_sort_config<256U>;
using type = merge_sort_config<limit_block_size<256U, sizeof(Key) + sizeof(Value)>::value>;
};

template<class Key>
struct merge_sort_config_900<Key, empty_type>
{
// static constexpr unsigned int item_scale =
// ::rocprim::detail::ceiling_div<unsigned int>(sizeof(Key), 16);

// using type = merge_sort_config<::rocprim::max(256U, 1024U / item_scale)>;
using type = merge_sort_config<256U>;
using type = merge_sort_config<limit_block_size<256U, sizeof(Key)>::value>;
};

template<unsigned int TargetArch, class Key, class Value>
Expand Down
16 changes: 14 additions & 2 deletions rocprim/include/rocprim/device/device_radix_sort_config.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -92,7 +92,13 @@ struct radix_sort_config_803
(sizeof(Key) == 8 && sizeof(Value) <= 8),
radix_sort_config<7, 6, scan, kernel_config<256, 13> >
>,
radix_sort_config<7, 6, scan, kernel_config<256, ::rocprim::max(1u, 15u / item_scale)> >
radix_sort_config<
6, 4, scan,
kernel_config<
limit_block_size<256U, sizeof(Value)>::value,
::rocprim::max(1u, 15u / item_scale)
>
>
>;
};

Expand Down Expand Up @@ -130,7 +136,13 @@ struct radix_sort_config_900
(sizeof(Key) == 8 && sizeof(Value) <= 8),
radix_sort_config<7, 6, scan, kernel_config<256, 15> >
>,
radix_sort_config<7, 6, scan, kernel_config<256, ::rocprim::max(1u, 15u / item_scale)> >
radix_sort_config<
6, 4, scan,
kernel_config<
limit_block_size<256U, sizeof(Value)>::value,
::rocprim::max(1u, 15u / item_scale)
>
>
>;
};

Expand Down