CUDA Dynamic Parallelism usage?

[tlemo]

What is the reason for using CUDA Dynamic Parallelism for the main simulation driver (calcSimulationTimestepKernel)?

alien/source/EngineGpuKernels/SimulationKernels.cuh

Lines 119 to 140 in 210fdd2

	__global__ void calcSimulationTimestepKernel(SimulationData data, SimulationResult result)
	{
	data.prepareForSimulation();
	result.resetStatistics();
	KERNEL_CALL_1_1(cudaApplyFlowFieldSettings, data);
	KERNEL_CALL(processingStep1, data);
	KERNEL_CALL(processingStep2, data);
	KERNEL_CALL(processingStep3, data);
	KERNEL_CALL(processingStep4, data, data.entities.tokenPointers.getNumEntries());
	KERNEL_CALL(processingStep5, data);
	KERNEL_CALL(processingStep6, data, result);
	KERNEL_CALL(processingStep7, data, data.entities.cellPointers.getNumEntries());
	KERNEL_CALL(processingStep8, data, result, data.entities.tokenPointers.getNumEntries());
	KERNEL_CALL(processingStep9, data);
	KERNEL_CALL(processingStep10, data);
	KERNEL_CALL(processingStep11, data);
	KERNEL_CALL(processingStep12, data, data.entities.particlePointers.getNumEntries());
	KERNEL_CALL(processingStep13, data);
	KERNEL_CALL_1_1(cleanupAfterSimulationKernel, data);
	result.setArrayResizeNeeded(data.shouldResize());

Also, somewhat related, every kernel launch is followed by a cudaDeviceSynchronize(). It seems an overkill in most places, am I wrong?

[chrxh]

During the calculations global synchronization points are needed, i.e., a kind of gpu-wide __syncthreads(). For example, while computing a new time step, first the forces acting on each particles are accumulated and in the next step the forces are applied.
In other cases, read and write operations are separated and performed in different processing steps to avoid conflicts. If one does not have a synchronization point in between, conflicts could easily arise. Regarding the processingSteps... in calcSimulationTimestepKernel, actually every cudaDeviceSynchronize() should be necessary.
But maybe there are nowadays more elegant ways for such global synchronization than to realize this via a kernel launch and cudaDeviceSynchronize()?

But I agree that at least for some "not very time-critical" kernels (e.g. editing operations) not every cudaDeviceSynchronize() is needed.
Because (partially for convenience) the general pattern here has become to start a single GPU thread first, which in turn invokes the required kernels between which global synchronization is needed.

[chrxh]

Thanks for the clarification! That helps me a lot since my understanding is quite incomplete.

I now think that I misunderstood the correct usage of cudaDeviceSynchronize.
In all these processing steps in the main pipeline, the only crucial thing is that the execution of the previous step must be completed. And that should be the case even without cudaDeviceSynchronize if I understand you correctly. So one can carefully remove most cudaDeviceSynchronize from the code.

I'm still curious to learn about the rationale behind this pattern. What kind of convenience does it provide? Transferring the kernel arguments instead of explicit data movement H->D?

Yes, actually that was one of my fear: More H->D copies and code from host side to reserve and free device memory for return values from kernels that need to pass to the next kernel.

[chrxh]

Here is an example of an iterative algorithm where I think it is more convenient to launch the kernel from the device code:

alien/source/EngineGpuKernels/EditKernels.cuh

Lines 664 to 667 in 210fdd2

	do {
	*result = 0;
	KERNEL_CALL(removeSelectedCellConnections, data, includeClusters, result);
	} while (1 == *result);

If this were written in host code, H -> D as well as D -> H transfers would be required. So I feel a bit indecisive if it would be worth the effort to completely remove the use of CDP. If the removal brings a performance gain (and for debugging reasons), one could just rewrite the main simulation pipeline (calcSimulationTimestepKernel) and leave other not so time-critical kernels as they are. How do you see this?

[tlemo]

Thanks. These iterative patterns present an interesting problem indeed. One upcoming complication is that the device-side cudaDeviceSynchronize is about to be deprecated starting in CUDA 11.6. So the CDP code might have to be rewritten anyway (without cudaDeviceSynchronize, one alternative would be to change the iteration to (tail) recursion)

I wouldn't dismiss the basic version, where the loop (and kernel launches for each iteration) are driven from the host.

[chrxh]

I've started to remove unnecessary cudaDeviceSynchronize in the time step calculation.
The removal of dynamic parallelism will be more time-consuming since I use a lot of iterations like above.

Such iterations are mainly used for 2 purposes:

1. To calculate connectivity components of a graph. The basic algorithm works as follows (I use indeed an optimized variant, but this is the basic idea):
   a) Kernel call: assign a unique number to each node (cell).
   b) Kernel call: Each node calculates the maximum from its own value and from that of the connected nodes, The result replaces the original value (with atomicExch). If the new value is different from the original one, a global variable repeat is set to true.
   c) Repeat b) as long as repeat = true.

In the end, two nodes are connected by a sequence of edges in the graph if and only if they have the same value. This algorithm can be parallelized very well.

2. In relatively rare cases it can happen that several threads have to process the same resource in certain operations. To avoid conflicts, in this case only one thread is allowed access. However, the thread that failed to gain access triggers a repeated kernel call.
   I know that if possible one should structure the data in such a way that such conflicts do not occur. In many cases I managed to do that. This approach is also only used for not so time-critical routines, i.e. not in the calculation of a simulation step, but in editing operations (e.g. if one wants to move or reassemble a subgraph (in alien a selected cell clusters) with the mouse).

[tlemo]

Thanks for explaining it. I was curious about the patterns which lead to CDP usage, and I think I understand it a bit better now.

[chrxh]

I've now removed the usage of CDP. The performance has improved a bit :) What I find quite bizarre though is (should be analyzed with a profiler):

alien/source/EngineGpuKernels/SimulationKernels.cu

Lines 108 to 114 in 4bb2783

	//This is the only kernel that uses dynamic parallelism.
	//When it is removed, performance drops by about 20% for unknown reasons.
	__global__ void nestedDummy() {}
	__global__ void dummy()
	{
	nestedDummy<<<1, 1>>>();
	}

[tlemo]

Great! How do you measure the performance? Do you have any external benchmarks?

What I find quite bizarre though is

I may be able to explain half of the mystery: CDP requires checkpoint injection (implicit sync points to allow things like device-side cudaDeviceSynchronize() to work). This machinery is only used if there's CDP usage anywhere in the program (determined statically, not at runtime). So the mere presence of nestedDummy() could introduce second order performance effects (as observed)

Why this translates to ~20% perf regression is a bit more strange. Like you said, it would require profiling to understand what's going on.

[chrxh]

That's interesting.
Yes, apparently the mere presence of the dummy "prevents" the performance degradation. I'd measured the number of time steps per second (without rendering the simulation). With the initial example the rate dropped from 380 to 310 and with larger simulation data from 54 to 40 timesteps/s.

I'll leave the dummy inside for now until the problem is analyzed with a profiler.