NVIDIA / NVIDIA/cccl

[BUG]: output iterators do not currently work with c.parallel merge_sort

Open
#3,722 0 comments 0 reactions 1 assignee Claimed by @NaderAlAwar View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

### Is this a duplicate?

- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this bug and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)

### Type of Bug

Runtime Error

### Component

cuda.parallel (Python)

### Describe the bug

Using output iterators with the c.parallel merge_sort results in incorrect output or memcheck errors. This test (which is currently commented out in #3636) reproduces this error:

```C++
struct random_access_iterator_state_t
{
int* d_input;
};

TEST_CASE("DeviceMergeSort::SortKeys works with output iterators", "[merge_sort]")
{
using TestType = int;
const int num_items = GENERATE_COPY(take(2, random(1, 1000000)), values({500, 1000000, 2000000}));

operation_t op = make_operation("op", get_merge_sort_op(get_type_info().type));
iterator_t output_keys_it =
make_iterator(
"struct random_access_iterator_state_t { int* d_input; };\n",
{"advance",
"extern \"C\" __device__ void advance(random_access_iterator_state_t* state, unsigned long long offset) {\n"
" state->d_input += offset;\n"
"}"},
{"dereference",
"extern \"C\" __device__ void dereference(random_access_iterator_state_t* state, int x) {\n"
" *state->d_input = x;\n"
"}"});
std::vector input_keys = make_shuffled_key_ranks_vector(num_items);
std::vector expected_keys = input_keys;

pointer_t input_keys_it(input_keys);
pointer_t input_items_it;
output_keys_it.state.d_input = input_keys_it.ptr;

merge_sort(input_keys_it, input_items_it, output_keys_it, input_items_it, num_items, op);

std::sort(expected_keys.begin(), expected_keys.end());
REQUIRE(expected_keys == std::vector(input_keys_it));
}

```

We get different behavior depending on the value of the input size `num_items`. When `num_items` is less than or equal to 1024, this works fine. For larger sizes, we either get incorrect output or `Invalid __shared__ read of size 4 bytes` in `DeviceMergeSortMergeKernel`.

I believe that the reason the issue occurs for sizes larger than 1024 is related to `num_passes` in `DispatchMergeSort`. For sizes large than 1024, `num_passes` is greater than 1.

Note that this issue only occurs when using an output iterator instead of a pointer.

The issue also occurs when we use an output iterator for items:

```C++
TEST_CASE("DeviceMergeSort::SortPairs works with output iterators for items", "[merge_sort]")
{
using TestType = int;
using item_t = int;
const int num_items = GENERATE_COPY(take(2, random(1, 1000000)), values({500, 1000000, 2000000}));

operation_t op = make_operation("op", get_merge_sort_op(get_type_info().type));
std::vector input_keys = make_shuffled_sequence(num_items);
std::vector input_items(num_items);
std::transform(input_keys.begin(), input_keys.end(), input_items.begin(), [](TestType key) {
return static_cast(key);
});
std::vector expected_keys = input_keys;
std::vector expected_items = input_items;

iterator_t output_items_it =
make_iterator(
"struct item_random_access_iterator_state_t { int* d_input; };\n",
{"advance",
"extern \"C\" __device__ void advance(item_random_access_iterator_state_t* state, unsigned long long
offset) "
"{\n"
" state->d_input += offset;\n"
"}"},
{"dereference",
"extern \"C\" __device__ void dereference(item_random_access_iterator_state_t* state, int x) {\n"
" *state->d_input = x;\n"
"}"});

pointer_t input_keys_it(input_keys);
pointer_t input_items_it(input_items);
output_items_it.state.d_input = input_items_it.ptr;

merge_sort(input_keys_it, input_items_it, input_keys_it, output_items_it, num_items, op);

std::sort(expected_keys.begin(), expected_keys.end());
std::sort(expected_items.begin(), expected_items.end());
REQUIRE(expected_keys == std::vector(input_keys_it));
REQUIRE(expected_items == std::vector(input_items_it));
}

```

### How to Reproduce

1. In `iterators.cpp` add back the changes made to the output iterator source in edd1e41dfab2b1f48019e0b285f4ce5798a49247
2. Add the above tests to `c/parallel/test/test_merge_sort.cpp` (and optionally comment out the other tests to avoid interference)
3. Compile and run using these commands
```bash
mkdir build & cd build

# build cccl.c with building tests enabled:
cmake .. -DCCCL_ENABLE_C=Yes -DCCCL_C_Parallel_ENABLE_TESTING=On -DCMAKE_BUILD_TYPE=Release

# build a test (as a side effect, builds cccl.c as well):
cmake --build . --target "cccl.c.parallel.test.merge_sort.cpp"

# run the test
./bin/cccl.c.parallel.test.merge_sort.cpp
```
4. (Optional) set `num_items` to 1025 to observe the incorrect output with no memcheck errors

### Expected behavior

The above test should pass but it is currently failing

### Reproduction link

_No response_

### Operating System

_No response_

### nvidia-smi output

_No response_

### NVCC version

_No response_

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.