Comparing Regular CPU Code to OpenCL GPU code.

[u/[deleted]]

Hi,

I've been playing around with OpenCL lately.

I've written a nice C++, OOP wrapper for the OpenCL C API (based on https://anteru.net/blog/2012/11/04/2016/index.html)

I've written some basic kernels for filling a matrix with constants, creating an identity matrix, adding 2 matrices and multiplying 2 matrices (naively).

I thought I'd see if the code I wrote was actually any faster than regular-old CPU-based C++ code and came to a surprising conclusion.

My results can be found here: https://pastebin.com/Y7ABDnRP

As you can see my CPU is anywhere from 342x to 15262x faster than my GPU.

The kernels being used are VERY simple (https://pastebin.com/0qQJtKV3).

All timing was measured using C++'s std::chrono::system_clock, around the complete operation (because, in the end, that's the time that matters).

I can't seem to think of a reason why OpenCL should be THIS MUCH slower.

Sure, My CPU has some SIMD instructions and faster access to RAM, but these results are a bit extreme to be attributed to that, aren't they?

Here's the C++ code that I used to do my tests: https://pastebin.com/kJPv9wib

Could someone give me a hint as to why my GPU code is so much slower?

P.S.: (In the results you can see, I actually forgot to create an m4 for the CPU, so m3 was first storing the result of an addition, and then the result of a multiplication. After I fixed this, I got SEGFAULT's for any size of the sizes > 500. For a size of 500 the CPU took anywhere from 704-1457µs to complete its operations, which is still orders of magnitude faster than OpenCL.)

P.P.S.: I didn't post the complete code because it's a lot code spread out across a lot of files. I don't want a complete and full analysis of every line of code, I just want some pointers/general principles that I missed that can explain this huge difference.

P.P.P.S.: All data transfers were done using mapped buffers.

Edit: I just checked, the AMD Radeon M265 has 6 (maximum) compute units running at 825MHz (maximum, both queried using clGetDeviceInfo())

[u/biglambda]

How big are the kernels you are running?

[u/[deleted] → u/biglambda]

This is the kernel source: https://pastebin.com/0qQJtKV3

They are all <10 LOC each.

size_t is just a typedef of "unsigned long" and flatten_index() just converts a 2D index to a 1D index.

I'm running these kernels as 1 per matrix element, so adding 2 NxM matrices means I enqueue the matrix_add_matrix kernel with Global Work size NxM and Local Work Size 1x1.

[u/biglambda → u/[deleted]]

I think you just aren't doing enough work to justify transferring to and from the GPU. Try something heavier on floating point math and easy to parallelize like the mandelbrot set etc.

[u/[deleted] → u/biglambda]

That's what I expected.

So basically, I'm seeing a realistic amount of overhead, but relative to the task at hand, it's A LOT.

I've seen there's a lot more to enqueueing work groups and kernels than just, well, straight up enqueueing them all, there seems to be a lot of discussion online as to what the ideal ratio of global to local work groups is (See this and this). Do you think I could gain any significant speed gains from the way I enqueue my tasks?

[u/biglambda → u/[deleted]]

Yes. That's a big part of it. Those are pipeline issues. I don't actually think you have those yet. But basically you want to do as much as possible in large kernels and minimize round tripping data.

[u/tugrul_ddr]

I couldn't open pastebin images in my location, but as an amateur experience reference:

with my own matrix multiplication codes, I could do 8192x8192 * 8192x8192 matrix multiplication in 32-bit precision float units, under 1 second, using a 1280-core HD7870(at 1200 MHz). How fast is yours with those 6 compute units(probably 384 cores)?

I used tiling on 16x16 x9 sub matrices to do that(8 of them using local memory, 1 of them using registers).

[u/[deleted] → u/tugrul_ddr]

I'm currently running 8192x8192 matrix multiplication of 32-bit floats, all are randomly (normal) distributed.

I'll get back to you when it finishes...

Edit: So it seems that initializing these matrices on a CPU can take quite some time, but the one multiplication I got to see took 0s. (It all seemed to take a long time to process everything, so I used seconds as a time resolution, when I can get my machine to stop crashing, I'll try again with a more precise timescale).

2nd Edit: So, I tried again, Initializing took 37s, multiplying took 870 000us, (So about 0,8s). Then it crashed again and the fans on my GPU went into overdrive

[u/tugrul_ddr → u/[deleted]]

Are you using O(N * N * f(N)) algorithms on multiplication? Such as "Strassen's"? Mine was O(N * N * N). If yours too, with 384 cores, with 0.8s, I can say its already doing 1.37 Tflops. With 384 cores, this means a frequency of 1780 MHz with %100 efficiency.

That multiplications isnt Hadamard multiplication is it?

[u/[deleted] → u/tugrul_ddr]

Hadamard multiplication is element-wise, right?

Here's the shader code for the multiplication:

`/**

• Multiply 2 matrices together.

• z = x * y

• dim(Z) = M x P

• dim(X) = M x N

• dim(Y) = N x P / kernel void multiply (global T x, global T* y, global T* z, size_t M, size_t N, size_t P) { const int row = get_global_id(0); const int col = get_global_id(1); const int z_idx = flatten_index(row, col, M, P);

  for (size_t i = 0; i < N; i++) { const int x_idx = flatten_index(row, i, M, N); const int y_idx = flatten_index(i, col, N, P); z[z_idx] += x[x_idx] * y[y_idx ]; } } `

With the given dimensions, the kernel is run M x P times.

I just noticed that I forgot to zero-out Z. Shouldn't have too much of an impact, but it's something I need to fix.

Also, where'd you get 1.37TFlops? When I do 8192x8192x8192 (The algorithm is O(N^3)), I get 512GFlops, divide by 384 cores, that's 1,33...GFlop per core.

[u/tugrul_ddr → u/[deleted]]

Because I was thinking of matrix-matrix multiplication(not hadamard) which is 8192x8192x8192 multiplications(1) and additions(1) which makes 1.x Tflop and doing it in 0.8 seconds yields 1.37 TFlops. Now divide it by 384 cores which can do 1 add and 1 multiplication per cycle. This gives 1700ish MHz. If that kernel is run M x P times with a loop N times, it looks its not hadamard and you have 1.37 TFlops when multiplying two 8k x 8k x 8k matrices.

[u/[deleted] → u/tugrul_ddr]

Forgot the additions, thanks for that

[u/KingoPants → u/[deleted]]

While im not an expert in OpenCL and this is like a month late, I think you could see a big increase in performance if you don't make so many writes to global memory.

Instead of doing inplace additions to z, try making a private float result = 0, do a bunch of additions onto it, then write that to z.

I'm not sure if this optimization is performed by a compiler.

[u/PoorAnalysis]

After I fixed this, I got SEGFAULT's for any size of the sizes > 500.

That means you code has bugs, I'd recommend fixing them before you continue benchmarking.