add tests for scan_by_key in-place execution

[gevtushenko]

Our scan by key implementation doesn't allow keys/result aliasing. This is caused by non-synchronized reading of d_keys_in[tile_base - 1] in the following code. Therefore, aliasing of keys and results introduces a data race. The following code can be used as a reproducer:

#include <thrust/scan.h>
#include <thrust/device_vector.h>

int main() {
  const std::size_t n = 64 * 1024 * 1024;
  thrust::device_vector<int> keys(n, 1);
  thrust::device_vector<int> values(n, 1);
  thrust::device_vector<int> output_storage(n);

  // thrust::device_vector<int> &output = values; // alias is OK
  thrust::device_vector<int> &output = keys; // alias is NOT OK

  thrust::exclusive_scan_by_key(keys.begin(), keys.end(), values.begin(), output.begin());

  if (output[n - 1] != n - 1) {
    std::cerr << "wrong result: " << output[n - 1] << " != " << n << std::endl;
  }
}

We also don't have tests for this use case. I can fix the code to guarantee this behaviour (just like in adjacent difference), but it'll slow down non-in-place execution in some cases. My suggestion is to have a quick fix of documentation.

[alliepiper]

Summarizing some offline discussions and testing:

• This has been broken for years, and is not a regression introduced by the recent changes to the scan_by_key implementation.
• The common case of overlapping values_in/values_out is fine, the only issue is overlapping keys_in/values_out, which is a somewhat odd usecase.
• Based on similar situations, fixing this is expected to introduce ~15% slowdown.

@senior-zero will be investigating this a bit more before we make a decision here.

[gevtushenko]

I've managed to fuse some kernels to reduce slowdown up to 6% on small to moderate problem sizes. The performance regression disappears at about 2^20 items for std::uint{16,32,64}_t and at about 2^22 for std::uint8_t.

[alliepiper]

6% for small inputs and no impact on larger inputs seems worthwhile to restore the documented guarantees. IMO we should update the implementation.

[gevtushenko]

@allisonvacanti I've bumped CUB version to one that incorporates updated scan by key. I'd like to re purpose this PR to update tests in thrust. Please, review the changes.

[gevtushenko]

run tests

[alliepiper]

Let's also add tests to the thrust/testing/scan_by_key.*.cu files to make sure the non-CUDA backends test this, too.

[gevtushenko]

run tests

[miscco]

Nitpick, no change requested: I would prefer if we put the whole test in its own block rather than only the check, as that helps readability a lot

Suggested change

	// in-place scans: keys/values aliasing
	thrust::inclusive_scan_by_key(h_keys.begin(), h_keys.end(), h_vals.begin(), h_output.begin());
	inclusive_scan_by_key_kernel<<<1,1>>>(exec, d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys.begin());
	{
	cudaError_t const err = cudaDeviceSynchronize();
	ASSERT_EQUAL(cudaSuccess, err);
	}
	ASSERT_EQUAL(d_keys, h_output);
	{ // in-place scans: keys/values aliasing
	thrust::inclusive_scan_by_key(h_keys.begin(), h_keys.end(), h_vals.begin(), h_output.begin());
	inclusive_scan_by_key_kernel<<<1,1>>>(exec, d_keys.begin(), d_keys.end(), d_vals.begin(), d_keys.begin());
	cudaError_t const err = cudaDeviceSynchronize();
	ASSERT_EQUAL(cudaSuccess, err);
	ASSERT_EQUAL(d_keys, h_output);
	}