Reduction Kernels locked to 1 block

[jplauzie]

Hi,

I was poking about the existing reduction kernels for #171 . It seems like in reduce.cu and cudalaunch.hpp, the reduction kernels seem to be hard locked to only use a single block (256 threads):

  dim3 gridDims(1);
  cudaStream_t s0 = getCudaStream();
  kernelfunction<<<gridDims, blockDims, 0, s0>>>(args...);

Is this intended/known bug? (Maybe left over from trying to debug the cudamemcpyasyncs?). This would be a pretty large performance penalty for reductions. For something like standardproblem3 or hysteresis which is all minimizes, it's something like ~45%, and probably scales with larger grid sizes. For any dynamic sims, the effect is probably smaller since the FFT kernels will dominate.

(There are some other potential optimizations in terms of things like warp shuffles as well, but that is probably better suited for a PR)

[JonathanMaes]

The intent of this is probably to make the reduction operation deterministic, unlike mumax³ where the usage of multiple blocks causes, for instance, inconsistent results for the minimizer (a fix for this will soon be pushed).

In mumax⁺, the minimizer does seem to be deterministic, so if we want to increase the number of blocks used, care must be taken to properly apply the reduction between the results of different blocks in order for the end result to remain deterministic (as will be done in mumax³).

[jplauzie]

There's a few things that can be done to make the minimizer more deterministic (in both mumax3 and mumax⁺), so I will wait to see what you've cooked up 😄 But yes, I do think it would be ideal to have multiple blocks if possible, since it will have a large impact on specifically stuff like hysteresis calculations.

Have you guys seen this (this might be what you're referring to)? Seems potentially very useful for determinism in floating point operations

https://developer.nvidia.com/blog/controlling-floating-point-determinism-in-nvidia-cccl