Add DeviceValue/HostValue

[sethrj]

@pcanal @tmdelellis I'm making a new branch (after merging physics-base) for the DeviceValue code because I think it'll need some iteration to get right. Right now with a simple struct, I don't think it's adding much value and it's adding a LOT of boilerplate code that clutters up the usage. Basically I have to add {} around everything to implicitly convert to a XValue, and add .value all over the place to get the pointed-to data. Because the struct allows implict construction and direct access of the data, I'm not sure it's doing much in that sense -- it's not like thrust::device_ptr which forces you to call device_ptr_cast wherever you use it.

Please take a quick look at some of the changes and you'll hopefully see what I mean about this first attempt being clumsy -- any suggestions on improvements?

[pcanal]

We could consider using just

Suggested change

	using DevicePointers = DeviceValue<span<byte>>;
	using HostPointers = HostValue<span<byte>>;
	using constDevicePointers = DeviceValue<span<const byte>>;
	using constHostPointers = HostValue<span<const byte>>;
	using HostPointers = span<byte>;
	using constHostPointers = span<const byte>;

unless the symmetry is more helpful.

[pcanal]

We might be able to avoid the double curly braces with

Suggested change

	T value;
	DeviceValue(T &&input) : value(std::move(input)) {};
	T value;

or

Suggested change

	T value;
	template <typename... ArgsTypes>
	DeviceValue(ArgsTypes... params) : value(params...) {};
	T value;

The later will make the compiler error message a bit more indrect when getting the parameter list wrong.

[pcanal]

and add .value all over the place to get the pointed-to data.

The only places that should be getting the pointed-to data are:

• the code that allocate or free the device memory.
• registeries that caches the address (but those could "just" keep the opaque-ish device pointer
• the kernel launchers

Another challenge maybe ParticleParamsPointers ParticleParams::device_pointers() const vs ParticleParamsPointers ParticleParams::host_pointers() const where we could/should have:

DeviceValue<ParticleParamsPointers> ParticleParams::device_pointers() const

because (as it is right now) device_pointers is a host function but its return value is usable only by a kernel launcher (and not by regular) code.
And if we go this direction it would seems that the line (inside ParticleParams::device_pointers) that reads:

result.defs = device_defs_.device_pointers().value;

to read

result = device_defs_.device_pointers();

but that would require a slightly different type, i.e. a version of ParticleParamsPointers that has a defs as a DeviceValue.

We could possibly achieve this will something like

template <typename T> using HostValue = T;

template <  template<typename> typename LocationWrapper = HostValue >
struct ParticleParamsPointers {
     LocationWrapper< span<ParticleTrackState> > defs_;
}

and then we could have:

ParticleParamsPointers<DeviceValue> ParticleParams::device_pointers() const

Underlying some of this rambling is that "DeviceValue/Pointer" is essentially a concept that is solely used in host code to manipulate memory allocate on device.

[pcanal]

Actually I am realizing than in the above I seem to also be thinking that rather that "host" the alternative to "device" is "local" (i.e. host in host code, device in device code) and local can just be straight pointers but at the interface (which is essentially only host code calling device code), there is a need for a "pointer in the other guy" concept (i.e. pointer value that can only be used in the interface routine (cudaFree, cudaCopy and kernel launch)

[pcanal]

And by kernel launch, we can either have:

__host___
void call_kernel( DeviceValue<Type> value, int blocksPerGrid, int threadsPerBlock, cudaStream_t stream) {
      some_kernel<<< blocksPerGrid, threadsPerBlock, 0, stream >>> ( value.value_ );
}
...
__global__
void some_kernel( Type value ) {
...

or

__global__
void some_kernel( DeviceValue<Type> in_value ) {
    Type value = in_value.value_;  // Or a pointer to Type or a reference to Type
    ....
}

where the former is likely the better one.

[sethrj]

@pcanal I let this drop in August because I got caught up in other stuff. I think you're right that the correct way to look at this is "device vs local" rather than "device vs host" -- I think I got buried in my original implementation here. I'm going to close this branch soon because I'm trying to clean up the upstream repo, but on Monday let's talk about device pointers and interfaces and nomenclature now that we have more use cases in the code than when we first started.