Clarify the residency of initial values for all Thrust/CUB APIs

[leofang]

xref: #2335 (comment)

some API does support initial value in device memory, but most (reduce included) do not.

[jrhemstad]

Can you be more specific? Each algorithm's API docs make it clear what the accessibility of the values are.

https://nvidia.github.io/cccl/cub/api/structcub_1_1DeviceReduce.html#_CPPv4I00000EN3cub12DeviceReduce6ReduceE11cudaError_tPvR6size_t14InputIteratorT15OutputIteratorT9NumItemsT12ReductionOpT1T12cudaStream_t

[leofang]

The quote was from @gevtushenko not me. Looking at the doc you link, I can't tell where init should reside (host or device or either).

[jrhemstad]

I understand Georgii's quote and I'm saying our docs already make it clear. For reduce, init is passed by value so it doesn't really make sense to talk about where it resides.

[leofang]

I am only more confused then. So if we say T init implies a host value should be used (for which I agree), are there APIs that allow a device value?

some API does support initial value in device memory

Perhaps the answer is simply: cub::Device always takes a host init value, and anything below (cub::Block etc) always takes a device init value? (Not sure about Thrust.)

[jrhemstad]

I think @gevtushenko was referring to the overload of DeviceScan that takes a FutureValue argument as the initial value which is unique in accepting a potentially device resident value.

https://nvidia.github.io/cccl/cub/api/structcub_1_1DeviceScan.html?highlight=future#_CPPv4I000000EN3cub10DeviceScan13ExclusiveScanE11cudaError_tPvR6size_t14InputIteratorT15OutputIteratorT7ScanOpT11FutureValueI10InitValueT14InitValueIterTE9NumItemsT12cudaStream_t