Simplify SiPixelFedCablingMapGPU SoA

[makortel]

I wanted to dump the SiPixelFedCablingMapGPU to a file for some standalone testing (we can talk about that next week), and easiest was to try out the suggestion #272 (comment) since the arrays were allocated to compile-time maximum size anyway.

[VinInn]

Looks definitively cleaner and simpler to manage (ok I am biased)

[fwyzard]

not relevant for this PR, but wouldn't it be simpler to use a "dumb" pointer (SiPixelFedCablingTree const *) instead of a const referent to a unique_ptr ?

[makortel]

I'd say it would be clearer to drop const& and take the unique_ptr by value as cablingMap.cablingTree() returns the unique_ptr by value

cmssw/CondFormats/SiPixelObjects/interface/SiPixelFedCablingMap.h

Line 35 in 30e75bb

std::unique_ptr<SiPixelFedCablingTree> cablingTree() const;

[fwyzard]

same for fedIds ?
cablingMap.fedIds() return an std::vector by value, so we could drop the const& there as well, and let the compiler move or even optimise it away

[fwyzard]

modToUnpDefault_ ?

[fwyzard]

size_ ?

[fwyzard]

last comment not relevant to this PR, but rather to CUDAESProduct in general :-)

The pattern seems to be:

• start with a class/struct for the actual data on th GPU

  Payload *cablingMapHost = nullptr; // pointer to struct in CPU

• define a wrapper

  struct PayloadWrapper {
    ~PayloadWrapper();
    Payload *payload = nullptr; // pointer to struct in GPU
  };

• add a CUDAESProduct data mamber:

  CUDAESProduct<PayloadWrapper> payload_;

• produce it for the gpu like this

Payload const* getGPUProductAsync(cuda::stream_t<>& cudaStream) const {
  const auto& data = payload_.dataForCurrentDeviceAsync(cudaStream, [this](PayloadWrapper& data, cuda::stream_t<>& stream) {
      // allocate
      cudaCheck(cudaMalloc(&data.payload, sizeof(Payload)));

      // transfer
      cudaCheck(cudaMemcpyAsync(data.payload, this->cablingMapHost, sizeof(Payload), cudaMemcpyDefault, stream.id()));
  });
  return data.payload;
}

Would it make sense to encapsulate more of the common part into CUDAESProduct ?

And/or to drop the PayloadWrapper in favour of a unique_ptr, possibly with a custom destructor ?

[makortel]

I'm all for encapsulating patterns, but I need to think this for a while. I made an issue #336 of it to remind.

[fwyzard]

I have not seen a measurable¹ impact on the throughput.

V100 on JetHT data

Average of 10 jobs running with 10 threads on a single GPU

reference :

1689.1 ±  24.8 ev/s
1674.2 ±  11.9 ev/s

#301:

1674.6 ±   7.6 ev/s

T4 on TTbar MC

Average of 10 jobs running with 8 threads on a single GPU

reference :

 720.9 ±   1.3 ev/s
 722.3 ±   1.1 ev/s

#301:

 719.9 ±   0.7 ev/s

¹ the measurements on the V100 seem to fluctuate a lot ...

[fwyzard]

is there a reason why at line 39 we use

for (unsigned int fed = startFed; fed <= endFed; fed++) {

instead of

for (unsigned int fed: fedIds) {

?

[makortel]

I don't know if it can happen, but if fedIds vector does not contain all the values between fedIds.front() and fedIds.back(), the result of those two is different.

[fwyzard]

Sure - I guess I was wondering if having all the values in between is a requirement or not.