A bunch of questions in cufinufft with respect to ntransf and streams

[chaithyagr]

I was curious about a bunch of questions (in cufinufft) concerning this repository that could help with performance:

1. When we launch NUFFT with ntransf>1, I see we launch all kernels on same stream on GPU in a loop. I was wondering why don't we launch them on a new stream so that they are executed in parallel, is it currently executed in parallel or am I missing something?

2. If we are anyway launching all kernels on default stream, they will run synchronously, hence I was wondering the purpose of cudaDeviceSynchronize here:

   finufft/src/cuda/3d/cufinufft3d.cu

   Lines 90 to 101 in 5c620c2

   	d_plan->c = d_cstart;
   	d_plan->fk = d_fkstart;
   	// Step 1: amplify Fourier coeffs fk and copy into upsampled array fw
   	cudeconvolve3d<T>(d_plan, blksize);
   	// Step 2: FFT
   	cudaDeviceSynchronize();
   	cufft_ex(d_plan->fftplan, d_plan->fw, d_plan->fw, d_plan->iflag);
   	// Step 3: deconvolve and shuffle
   	ier = cuinterp3d<T>(d_plan, blksize);

   
   This can add minor latencies if it isn't needed.

3. Finally I was wondering if we can add API to potentially send a stream pointer so that everything is launched on this stream? (Or a list of streams in case of ntransf>1) . This can really help a lot in building on concurrency with a IO task happening in background for next k-space data, potentially giving great speedups. I can work on a PR for this if required.

[ahbarnett]

Dear Chaithya,
These interesting questions on the cuda details are a bit beyond my expertise, but I hope Robert @blackwer,
Joakim @janden, and the original author Melody @MelodyShih may be able to add insight - they are the kind of thing (relating also to using multiple GPUs at once) they have been discussing.
I should also ping Marco @DiamonDinoia another contributor and Johannes @JBlaschke on the paper and using for multi-GPU.
I suggest the device synchronize (which is only called once) is related to calling cuFFT.
Thanks all, Alex

[blackwer]

1. When we launch NUFFT with `ntransf>1`, I see we launch all kernels on same stream on GPU in a loop. I was wondering why don't we launch them on a new stream so that they are executed in parallel, is it currently executed in parallel or am I missing something?

Launching on multiple streams is generally helpful for overlapping host <==> device transfers with kernel execution, which is on the agenda of improvements to consider. It can be used for concurrent kernel execution, but, unless you're working with very small transforms (where you probably shouldn't be using cufinufft, or even finufft for...) it's unlikely that concurrent kernel execution will be particularly helpful. In most cases the GPU occupancy should be high enough that this isn't a problem.

Regardless, we'll be looking into general optimizations in the near future, and I'll investigate how to best utilize streams.

2. If we are anyway launching all kernels on default stream, they will run synchronously, hence I was wondering the purpose of `cudaDeviceSynchronize` here: https://github.com/flatironinstitute/finufft/blob/5c620c2a9068bdc20a3ad4b0b9cf27fea9c390c6/src/cuda/3d/cufinufft3d.cu#L90-L101
   
   This can add minor latencies if it isn't needed.

You're right. This should be a pedantic call and unnecessary. The cost should be order nanoseconds iirc though. I'll remove it on my optimization pass.

3. Finally I was wondering if we can add API to potentially send a stream pointer so that everything is launched on this stream? (Or a list of streams in case of `ntransf>1`) . This can really help a lot in building on concurrency with a IO task happening in background for next k-space data, potentially giving great speedups. I can work on a PR for this if required.

This will require some thought. By IO tasks do you mean moving data to the GPU, or loading it from some data store? In the latter case, you can thread that off yourself straightforwardly right? In the former, I'm skeptical this will help more than a few percent, but open to have my mind changed. Transfers are typically a tiny tiny fraction of the computation time. Do you have some rough numbers for timings and/or the type of data you're working with? Is the GPU reporting any kind of underutilization? Are you using the python or C/C++ API?

[blackwer]

@blackwer Here is an interesting comprehensive comparision with gpuNUFFT and for 3D, we can clearly see gpuNUFFT outperforms as it uses concurrency and in general I think its spread and interp is faster if I am right. sadly it isn't maintained and needs some major fixes (which I will get at if speed gets crucial)

No link/attachment. Can you resend?

It is indeed between transfer between CPU and GPU. Depending on the size of matrices we can expect maybe 10-20% from what I remember. I can setup a perf test and send the nsys outputs if it helps.

Roughly how big are the transforms and how many are done in a single nufft call? Are they type 1 or 2? Is calling cufinufft the bulk of your program's execution time? Regardless, if you have some synthetic driver code/tests/perf graphs all the better, but it's certainly not necessary. I'm already working on some synthetic driver code already for performance benchmarking.

[chaithyagr]

No link/attachment. Can you resend?

Ouch my bad: flatironinstitute/cufinufft#126 (comment)

[chaithyagr]

Roughly how big are the transforms and how many are done in a single nufft call? Are they type 1 or 2? Is calling cufinufft the bulk of your program's execution time? Regardless, if you have some synthetic driver code/tests/perf graphs all the better, but it's certainly not necessary. I'm already working on some synthetic driver code already for performance benchmarking.

I guess most answers in that comment

[chaithyagr]

Adding here a recent pic for 3D matrix size of (128, 192, 192). As you see the memory copy is a significant overhead!

[blackwer]

Are those memory operations yours, or ours? I.e. Is this a single call with a large ntransf, or are those your host-to-device transfers? This looks like a single transform and then a host->device, then a single transform, etc, rather than batching. Right?

I tried speeding up things internally with streams (each stream gets maxbatchsize=8 transforms), and running with more than one stream always makes things slower (down by a factor of 30-50% in most cases I tried). You gain very very marginal kernel concurrency (on the tails of the kernels), but the kernels execute slower on average. I haven't quite figured out what's going on, but I suspect it's a caching issue. Most of cufinufft is memory bound, so swapping out the data you're operating on can only hurt you. Just a theory though.

That said...

The host <==> device problem is very real and probably a much more fruitful target. Type 1 and 2 transforms below. Amortized cost is the entire cost of the run, from start to finish. GPU is an A6000, though the qualitative results are similar on the v100/a100/h100. Note this is for 32 transforms at once. The net speed gets worse without batching.

# prec = f
# type = 1
# n_runs = 1
# N1 = 128
# N2 = 128
# N3 = 160
# M = 6136781
# ntransf = 32
# method = 1
# kerevalmethod = 1
# sort = 1
# tol = 1e-05

event	count	tot(ms)	mean(ms)	std(ms)	nupts/s	ns/nupt
host_to_device	1	420.590179	420.590179	0.000000	0.0	0.0
makeplan	1	1.771328	1.771328	0.000000	0.0	0.0
setpts	1	1.415168	1.415168	0.000000	1.38766e+11	0.007206
execute	1	1815.781372	1815.781372	0.000000	1.0815e+08	9.246406
device_to_host	1	34.795902	34.795902	0.000000	0.0	0.0
amortized	1	2274.374512	2274.374512	0.000000	8.63433e+07	11.581675

# type = 2

event	count	tot(ms)	mean(ms)	std(ms)	nupts/s	ns/nupt
host_to_device	1	341.503967	341.503967	0.000000	0.0	0.0
makeplan	1	10.117152	10.117152	0.000000	0.0	0.0
setpts	1	1.699840	1.699840	0.000000	1.15527e+11	0.008656
execute	1	285.270020	285.270020	0.000000	6.8839e+08	1.452665
device_to_host	1	172.218918	172.218918	0.000000	0.0	0.0
amortized	1	810.825745	810.825745	0.000000	2.42194e+08	4.128924

I have some ideas for how to deal with this, given the performance issues with running execute on multiple streams simultaneously. I'll ping you when/if I have something to test.