Optimise OpenCL kernels

[hjabird]

Surely the code can be faster?

[hjabird]

OpenCL particle induced velocity:

See commit ae683f3f51ec4d2f70556ae6d08d4a6dc0561543

Approximately 7.5% faster:

• Move 1/4pi term to host
• Precompute 1 / regularisation_radius
• Winckelmans kernel: g = ... / powr(rho, 2.5) -> ... * rsqrt(pown(rho, 5))
• Neaten with radd = length(rad)
• Remove if(all(isequal(mespnt, particlepnt))) and replace with isnormal(..) ? ret : 0.0 0.0
• Use knowledge that get_global_id(0) == get_local_id(0)

Not used:

• fast_length instead of length results in ~7% faster, but this results in far lower precision.

Verified good:

• Workgroup size seems good.
• Reduction method seems good. I tried a scheme where reductions was performed in registers rather than local data store. This slowed down GPU considerably.
• Host memory prep only uses a few % of runtime.

To do:

• GPU loop unrolling?

NB:

• Benchmarking performed using a Julia script. The times are for the ind_vel only. The julia script adds overhead, dominating the low end results and adding a couple of seconds at around a million particles.

[hjabird]

commit d1a77dcc13d4869931b6ec830c926f492fb796ca

ind_dvort is approximately 27% faster.

• Move 1/(4 pi reg_rad^3) to host
• Precompute 1 / reg_rad
• Winckelmans kernel: same rsqrt(pown(..)) trick as above for both f and g terms.
• Same removal of if with ret = isnormal(ret) ? ret : (float3)(0.f, 0.f, 0.f) as above

As above, there is some overhead in the Julia script used for benchmarking. 3% in worst case.

[hjabird]

commit 6d962b76d15e95af67105c30e920fa670d2d24f1

ind_dvort is 1.5% faster

• Use knowledge that get_global_id(0) == get_local_id(0)
• use repetition of 1/(rho^3)

[hjabird]

commit 4db32f2e79e5eb11dea8f06539bd7f20bdaa5337

Changing to transpose method for dvort calculations allows reusing a cross product. 7% faster.