NVIDIA / NVIDIA/cccl

Merge sort key type selection

Open
#903 2 comments 0 reactions 1 assignee Claimed by @wmaxey View on GitHub
cub
Dominant language
C++
Stars
2.5k
Forks
487
Avg merge
2d 7h
Merged PRs (30d)
296

Description

Before porting to CUB, Thrust implementation of merge sort didn't use to have `*copy` version. When introducing `Copy` overload, I followed the CUB generic scheme of selecting output iterator value type. After that, we changed type selection for most algorithms so it no longer follows these scheme. Merge sort implementation, however, wasn't changed. This leads to compilation failures in some cases. @harrism provided the following [reproducer](https://godbolt.org/z/jhn5xeG1P):
```cuda
#include
#include

#include

#include

template
struct trajectory_comparator {
__device__ bool operator()(Tuple const& lhs, Tuple const& rhs)
{
auto lhs_id = thrust::get<0>(lhs);
auto rhs_id = thrust::get<0>(rhs);
auto lhs_ts = thrust::get<1>(lhs);
auto rhs_ts = thrust::get<1>(rhs);
return (lhs_id < rhs_id) || ((lhs_id == rhs_id) && (lhs_ts < rhs_ts));
};
};

struct foo {
float x;
float y;
};

struct bar {
float x;
float y;
};

struct foo_to_bar {
bar operator()(foo const& f) { return bar{f.x, f.y}; }
};

struct double_foo {
foo operator()(foo const& f) { return foo{f.x * 2.0f, f.y * 2.0f}; }
};

int main(void)
{
thrust::device_vector keys(1000);
thrust::sequence(keys.begin(), keys.end());

thrust::device_vector values(1000);

thrust::device_vector keys_out(1000, 0);

#if 0 // transform iterator input and output types the same
thrust::device_vector foo_values_out(1000);
auto values_transformed_out = thrust::make_transform_output_iterator(foo_values_out.begin(), double_foo{});
#else // transform iterator input and output types different
thrust::device_vector bar_values_out(1000);
auto values_transformed_out = thrust::make_transform_output_iterator(bar_values_out.begin(), foo_to_bar{});
#endif

std::size_t temp_storage_bytes = 0;
cub::DeviceMergeSort::SortPairsCopy(nullptr,
temp_storage_bytes,
keys.begin(),
values.begin(),
keys_out.begin(),
values_transformed_out,
1000,
thrust::less{});

void* temp_storage = nullptr;
cudaMalloc(&temp_storage, temp_storage_bytes);

cub::DeviceMergeSort::SortPairsCopy(temp_storage,
temp_storage_bytes,
keys.begin(),
values.begin(),
keys_out.begin(),
values_transformed_out,
1000,
thrust::less{});

return 0;
}
```
The issue is related to the fact that we use output iterator value type to instantiate block load facility and then provide input iterator values that are not accepted by the block load member functions. We have to investigate if changing key/value type selection to be based on the input iterators is possible.

Contributor guide

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.