Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Add example that profiles parallel sum #774

Open
wants to merge 5 commits into
base: main
Choose a base branch
from
Open
Changes from 1 commit
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
66 changes: 66 additions & 0 deletions examples/demo_flops.py
Copy link
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Could you look over the CI failures?

Original file line number Diff line number Diff line change
@@ -0,0 +1,66 @@
import pyopencl as cl
import numpy as np
import matplotlib.pyplot as plt


src = """
__kernel void sum(__global T *x, __global T *y, __global T *z) {
const int i = get_global_id(0);

z[i] = x[i] + y[i];
}
"""


# allocates buffers of increasing size, for each run do a parallel sum interpreting
# the buffer as an array of i8, i16, ...
# profile the kernels to find the throughput in GFLOPS, useful to estimate the raw computational speed of the hardware
if __name__ == '__main__':
types = [
('i8' , 'char' , 1),
('i16', 'short' , 2),
('i32', 'int' , 4),
('i64', 'long' , 8),
# ('f16', 'half' , 2),
('f32', 'float' , 4),
('f64', 'double', 8)
]


ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx, properties = cl.command_queue_properties.PROFILING_ENABLE)


buffer_size = [2 ** i for i in range(10, 31)]
data = np.zeros((len(buffer_size), len(types)))

for row, nbytes in enumerate(buffer_size):
x = cl.Buffer(ctx, cl.mem_flags.READ_ONLY, nbytes)
y = cl.Buffer(ctx, cl.mem_flags.READ_ONLY, nbytes)
z = cl.Buffer(ctx, cl.mem_flags.WRITE_ONLY, nbytes)

for col, (label, literal, sizeof) in enumerate(types):
sums = nbytes // sizeof
header = f'#define T {literal}\n'
kernel = cl.Program(ctx, header + src).build().sum

event = kernel(queue, (sums,), None, x, y, z)
FattiMei marked this conversation as resolved.
Show resolved Hide resolved
Copy link
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

It's generally good practice to do a few "warmup" rounds before timing, to better measure the steady-state rate.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

There are problem with caches however. In cpu runs I get crazy GFLOPS for medium size arrays because they already live in the cache, gpu doesn't seem to suffer from this problem.
But with the new commits one could decide to do no warmup runs and only one hot run so it's ok

event.wait()

FLOPS = 1e9 * sums / (event.profile.end - event.profile.start)
GFLOPS = FLOPS / 1e6

data[row, col] = GFLOPS
Copy link
Owner

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Arguably, this workload will be bandwidth-bound, so GB/s will be the more appropriate measure.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This decision was made because it's common to evaluate gpu performance based on TFLOPS (and this number is computed with similar workloads) and especially highlights the fact that of course the flops go up when working with smaller types


x.release()
y.release()
z.release()

for col, (_, label, _) in enumerate(types):
plt.semilogx(buffer_size, data[:, col], label = label)

plt.title(f'{ctx.devices[0].name}')
plt.legend()
plt.xlabel('sizeof(vector)')
plt.ylabel('GFLOPS')
plt.show()
Loading