Update DeviceScan to pass Thrust's scan tests

[alliepiper]

Thrust will be switching to cub::DeviceScan to replace its custom scan
implementation. This patch addresses some issues found by the Thrust
tests:

• Initialize unused BlockLoad items with values known to be in the input
  set. This fixes the TestInclusiveScanWithIndirection Thrust test by
  keeping the plus_mod3 functor indices valid.
• Use OffsetT instead of int to hold indicies in AgentScan. This
  fixes the Test*ScanWithBigIndexes Thrust tests by not truncating
  the input problem size.
• Use BLOCK_[STORE|LOAD]_WARP_TRANSPOSED_TIMESLICED instead of
  BLOCK_[STORE|LOAD]_WARP_TRANSPOSED when the intermediate type is
  larger than 128 bytes. This keeps shared memory buffers from growing
  too large in the TestScanWithLargeTypes Thrust test.

[alliepiper]

Requires NVIDIA/thrust#1304

[alliepiper]

DVS CL 29194365

[RAMitchell]

Not sure if this is covered already by this PR, but discard iterators typicaly don't work with cub interfaces because they can have value_type=void which propagates to OutputT, which cub tries to instantiate. I have been working around this as follows:

// Change the value type of thrust discard iterator so we can use it with cub
template <typename T>
class TypedDiscard : public thrust::discard_iterator<T> {
public:
  using value_type = T;  // NOLINT
};

[alliepiper]

@RAMitchell As of 2fba463, DeviceScan uses the input iterator's type for the scan accumulator, rather than the output iterator's type. But even before that, it had explicit checks for output type == void, so it appears that this was fixed at some point earlier.

Let us know in a new issue if you notice this happening in other algorithms.

[alliepiper]

DVS CL: 29264972