Port adjacent difference into CUB

[gevtushenko]

The PR contains:

1. a port of thrust::adjacent_difference algorithm;
2. deprecation of FlagHeads and FlagTails methods from BlockAdjacentDifference structure;
3. fixed API for BlockAdjacentDifference along with documentation and tests.

Note that the PR is based on some of the features introduced in MergeSort porting.

[alliepiper]

This is a great first pass 👍 Let me know when you're ready for a re-review.

[alliepiper]

LGTM -- just some comments on documentation, otherwise this is ready to start testing 👍

[alliepiper]

(continuing from #331 (comment))

Gotcha, so the last 5 values are unchanged random values that are ignored by the algorithm.

There's another convention that uses - to stand in for such values, e.g.

cub/cub/device/device_histogram.cuh

Lines 181 to 182 in d6ba3cf

	* float* d_samples; // e.g., [2.2, 6.1, 7.1, 2.9, 3.5, -, -,
	* // 0.3, 2.9, 2.1, 6.1, 999.5, -, -]

I think it's a little easier to parse with the non-numeric characters, since it's immediately clear that they can't be involved in the calculation. If you agree, let's change this to

* `{ ...3], [4,2,1,1], [1,1,1,1], [2,3,3,-], [-,-,-,-] }`.
* `{ ..., [-1,2,1,0], [0,0,0,-1], [-1,0,3,-], [-,-,-,-] }`.

[gevtushenko]

Gotcha, so the last 5 values are unchanged random values that are ignored by the algorithm.

These values aren't ignored by the algorithm:

As the documentation states, the input value is copied without modifications if the neighbour value is out of bounds.

If input and output parameters point to different memory locations, these values will be copied without modification. I'm afraid that non-numeric characters would mean that the output values are unchanged at all, for example:

input { 1, 2, 3 };
output { 4, 5, 6 };
valid_items = 0;

adjacent_difference(input, output, valid_items);

// output should be { 1, 2, 3 } and not { 4, 5, 6 }

The same with std::adjacent_difference documentation. The output for:

std::vector v {2, 4, 6, 8, 10, 12, 14, 16, 18, 20};

is

2 2 2 2 2 2 2 2 2 2

rather than:

- 2 2 2 2 2 2 2 2 2

Do you think it'll be clear from the - notation that the values are copied?

[alliepiper]

That's a good point -- I'm fine leaving this as-is.

[alliepiper]

s/on/of/

(this is repeated in a few other docstrings)

[gevtushenko]

There are quite a few places with this typo:

rg "threads on type" | wc -l
33

I'll fix this in the block_scan, block_discontinuity and block_reduce documentation as well.

[gevtushenko]

gpuCI: NVIDIA/thrust/pull/1577
DVS: 30766806