Audio buffer processing - Rapidly enqueue thousands of kernel runs sequentially on GPU then complete? Vs. just using `run()` in a loop (which takes too long)?

[jonmdev]

I am working on applying the OpenCL Wrapper (which is fantastic by the way and thanks) to audio processing.

I have written kernel scripts which are highly parallelized (using Memory<float>) but kernel operations must be run sequentially many times to process an audio buffer block.

Simulation of Processing Audio Buffer

Audio typically processes at 44100 Hz with an acceptable upper buffer size limit of 1024. This means we have ~23 ms to process a buffer (~42 buffers per second ie. ~42 "fps")

Hypothetically, for simplicity sake, let's say for this project here we have in our kernel.cpp these needed functions:

kernel void add_kernel(global float* A, global float* B, global float* C) { 
	const uint n = get_global_id(0);
	C[n] = A[n]+B[n];
}

kernel void smooth_kernel(global float* C) {
	const uint n = get_global_id(0);
	const uint N = get_global_size(0);

	if (n > 0 && n < N - 1) {
		C[n] = (C[n] + C[n] + C[n - 1] + C[n + 1]) * 0.25f;
	}
}
kernel void do_nothing_kernel(global float* C) {
}

So say for example, as we see here, add_kernel and smooth_kernel are both running parallelized operations, but cannot be combined into one kernel as the smoothing depends on the adding being done already.

Then in main.cpp, we have to replace your code:

#include "opencl.hpp"
//===============
//TIMER
//===============
//GET TIMER START TIME 
inline std::chrono::high_resolution_clock::time_point getStartTime() {
	std::chrono::high_resolution_clock::time_point start(std::chrono::high_resolution_clock::now());
	return start;
}
//GET TIMER FINISH TIME
inline double getTimeEllapsedMs(std::chrono::high_resolution_clock::time_point start) {
	std::chrono::duration<double> duration = std::chrono::high_resolution_clock::now() - start;
	double stopwatchMs = duration.count() * 1000;
	return stopwatchMs;
}

//=============
//MAIN
//=============
int main() {
	Device device(select_device_with_most_flops());

	const uint N = 1024u;
	Memory<float> A(device, N);
	Memory<float> B(device, N);
	Memory<float> C(device, N);

	for(uint n = 0u; n < N; n++) {
		A[n] = 3.0f;
		B[n] = 2.0f;
		C[n] = 1.0f;
	}
	A.write_to_device(); 
	B.write_to_device();
	C.write_to_device();

	print_info("Value before kernel execution: C[0] = " + cl_to_string(C[0]));

	//define kernels:
	Kernel add_kernel(device, N, "add_kernel", A, B, C);
	Kernel smooth_kernel(device, N, "smooth_kernel", C);
	Kernel do_nothing_kernel(device, N, "do_nothing_kernel", C);

	//===============================================
	//LOOP TEST TO SIMULATE PROCESSING AUDIO BUFFER
	//===============================================
	auto startTime = getStartTime();
	int numSamples = 1024;
	for (int i = 0; i < numSamples; i++) {
		add_kernel.run();
		smooth_kernel.run();
		do_nothing_kernel.run();
	}
	int timeMs = getTimeEllapsedMs(startTime);

	//===============================================
	//now hypothetically read output from processing
	C.read_from_device(); 

	//print result
	print_info("Value after kernel execution: C[0] = " + cl_to_string(C[0]) + " C[] " + cl_to_string(C[5]));
	print_info("Time Taken: " + cl_to_string(timeMs) + " ms");

	//takes 80-100 ms to complete, which is too long

	wait();
	return 0;
}

Result

I am getting with this operation 80-100 ms roughly to complete. This is far too long. This is matching my performance in my real world project where basic kernel runs are adding up to 40-120 ms over the course of an audio buffer.

As noted, we have only 23 ms to complete the buffer.

I note the majority of the time is just spent on running the kernels (ie. even do_nothing_kernel takes a very long time). This is presumably scheduling and thread/event coordination etc.

Idea

In this hypothetical, if we could move the sample loop audio buffer iteration to the kernel, so we only have to run a process_full_buffer kernel once I presume this would be faster.

But within that we would still need sequential processing on the GPU. Ie. The smooth_kernel function cannot run until the add_kernel is done on each sample. The smooth and add functions are parallelized but they must be run sequentially, and each audio sample must also be processed sequentially.

Is there any method other than run() or way you can imagine to enqueue and run say 1024 sequential sets of 3-5 kernel operations in one blast? ie. Send thousands of kernel operations for the GPU to coordinate and complete on its own in one batch, then let the thread continue when that is finished?

ie. So it won't take 80-100 ms just to coordinate or work through it all?

I presume it is all the "start run", "finish run", "trigger event that is done" etc that is slowing things down. If I could just hand the GPU as a batch 1024-4096 kernel operations and say "do those in order and tell the program when you're done at the end so it can continue" this would probably solve the problem. But I don't know how.

Like if we could make a std::vector<Kernel> kernels, add 1024-4096 kernels to it, then run an operation to pass all these to the GPU to do in order and then have the CPU thread continue after it concludes all of them as a batch.

Is such a thing possible? If so, how can you imagine it?

Thanks

Thanks for the great project and any ideas. I don't know any OpenCL. I picked your project up and within 2-3 days of porting it over to my actual projects I can run GPU code. So this is a testament to your great design. I just hope there is some way around this type of lag or I am not sure what to do.

(Minor Suggestion)

(Note that I renamed all your to_string function to cl_to_string in my code above as it was causing confusion otherwise when I brought over to projects that have std::to_string. Unrelated, but I think this would improvement for the project in any case. Perhaps you may wish to consider it.)

[ProjectPhysX]

Hi @jonmdev,

the kernel.run() command puts the kernel in the GPU queue and then waits (blocks the CPU thread) until the GPU sends back the signal that execution of the queue has completed. This is fine for kernels with rather long execution time, much larger than the latency to dispatch the kernel.

In your case dispatch latency is much more important. The solution is simple:

smooth_kernel.run(numSamples);

This will put the kernel numSamples times in the queue and then only once at the end wait for the full queue to complete.

You can also only enqueue with smooth_kernel.enqueue_run(numSamples) and manually wait for the queue to finish with smooth_kernel.finish_queue(). This can be useful if you want to fill the queue with several different kernels and only finish once in the end.
But note that you need a finish_queue() every once in a while, otherwise the queue blows up when you just add millions of kernels to it.

To the to_string issue: I've written my own to_string conversion functions because the default ones in std:: library suck; they don't convert with full precision, which is bad design and leads to many problems down the line when for example embedding floating-point constants in OpenCL C code.
You are likely facing a collision because somewhere you have a

using namespace std;

It's better to not force an entire namespace with a header on all the code below, but only use it for the functions/structs that you really need. For example:

using std::string;

Kind regards,
Moritz.

[jonmdev]

Thank you Moritz. I did some experiments and this cuts down the processing time considerably. However, I would still appreciate your further thoughts on it based on the new data I have accumulated.

1) KERNEL QUEUING EFFICIENCY ISSUES

In these cases, the functions are negligible as they are very simple. What we can see even with the following completely empty Kernel code:

kernel void do_nothing_kernel(global float* C) {
}

When used as:

int numSamples = 1024;
for (int i = 0; i < numSamples; i++) {
	do_nothing_kernel.enqueue_run();
	do_nothing_kernel.enqueue_run();
	do_nothing_kernel.enqueue_run();
	do_nothing_kernel.enqueue_run();
}
do_nothing_kernel.finish_queue();

Every extra queued kernel adds considerably to the total processing time. With one do_nothing_kernel run in the loop, I get 2 ms run time for the whole "audio buffer." With 4 in the loop, I get 10-12 ms.

Ie. Even though there is nothing to process, and we are now queuing, the time can add up quite considerably. This is relevant, eg. in Audio, because as noted, we only have 12-24 ms per audio buffer. And this allowed time is not just for one process, but ALL processes (other audio processes etc.).

All the time is spent on the queueing. Commenting out the finish_queue command makes no difference.

It seems that if we have to run 4+ Kernels to process the data stepwise per sample (due to functions like the smoothing above that must be done after prior step finishes), we are already eating up a massive amount of total available processing time just dealing with the overhead of qeueing the functions. This is without even a single useful calculation done yet.

Based on your understanding of OpenCL and your wrapper, is this inherent to OpenCL? Or is it perhaps specific to the architecture of the wrapper you built? ie. Could you imagine any solution to this, or is this just how it needs to be?

I have zero experience with this (just started "GPU coding" in the past few days based on your project) so I am not sure. If this is inherent to OpenCL, do you know if CUDA might offer any technology that can circumvent this?

12 ms to just enqueue sufficient (empty) functions to process an audio block is kind of a deal killer for anything audio related, so I am curious if you have any ideas at all. It is important to have some idea if this is perhaps solvable even theoretically. Thanks for any thoughts.

2) IS FINISH_QUEUE A "STATIC"/"GLOBAL" FUNCTION?

I note that it does not seem to matter which kernel I run finish_queue() on. Is that correct? Ie. If there are four kernels being queued repeatedly (or even one that is not), it makes no difference which runs finish_queue().

Ie. As long as it is run on a Kernel that exists, the OpenCL GPU queue is global, and this forces all of it to run through no different depending on which object it is run on. Seems this way, would just like to confirm.

3) DEBUGGING KERNEL CODE

Last question: Is there any way to debug anything out of the Kernel code? I have found it hard to figure out what is actually happening inside the Kernels. Is there any way to speak out of the Kernel or even write a log to disk or something?

Thanks again for your great project and quick and thorough reply. 🙂

[jonmdev]

I am sorry @ProjectPhysX , but I believe I may have misunderstood something about how kernels function.

My problem can be solved by one thnig if this is true. If you write a kernel like (again just simple hypothetical, not useful):

kernel void processAudioBlock(global float* A, global float* B, global float* C) { 

	const uint n = get_global_id(0);
	const uint N = get_global_size(0);

	//iterate through 1024 times for 1024 samples in audio buffer
	for (int i=0; i< 1024;i++){

		//Step 1: perform function that does not depend on neighbors
		C[n] = A[n]+B[n];

		//Step 2: perform function that DOES depend on neighbors (and thus needs Step 1 complete for all other samples)
		if (n > 0 && n < N - 1) {
			C[n] = (C[n] + C[n] + C[n - 1] + C[n + 1]) * 0.25f;
		}
	}
}

Does step 1 always complete for all GPU cores simultaneously before step 2 is run? Ie. is it valid to write it this way, or will each GPU core be potentially desynchronized?

I was under the assumption each core may run on its own time frame and thus you needed separate kernel calls to keep each step synchronized. If each core is always synchronized at each step of the kernel function then I can easily run the whole block in one kernel and this solves everything.

This would be very very good. :)

Still curious about the Debugging question if you have any solution and if I am correct if running the finish function is global and kernel independent also either way.

[jonmdev]

Update:

I have done some playing around and further reading. CUDA provides the same lag in queuing massive amounts of sequential kernels (actually slightly longer time than with the OpenCL wrapper).

The best solution for me seems to be doing CUDA with maximum 32 threads as this allows perfect synchronization as needed within one kernel with the __syncwarp() function.

Then I can run the whole buffer inside one kernel and just run __syncwarp() at various stages to resync.

So that is what I will do.

Again, thanks for your project in any case. It has been very helpful to me to get into GPU programming. I even copied your Memory style of memory management into a similar class for CUDA memory management as it makes the whole thing so much slicker with that approach.

So thanks for the great project and inspiration. All the best.🙂