-
Notifications
You must be signed in to change notification settings - Fork 245
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
base: main
Are you sure you want to change the base?
Changes from 1 commit
15d981a
dfef25a
df99f35
72b1ab4
766873e
File filter
Filter by extension
Conversations
Jump to
Diff view
Diff view
There are no files selected for viewing
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
There was a problem hiding this comment. Choose a reason for hiding this commentThe 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. There was a problem hiding this comment. Choose a reason for hiding this commentThe 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. |
||
event.wait() | ||
|
||
FLOPS = 1e9 * sums / (event.profile.end - event.profile.start) | ||
GFLOPS = FLOPS / 1e6 | ||
|
||
data[row, col] = GFLOPS | ||
There was a problem hiding this comment. Choose a reason for hiding this commentThe 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. There was a problem hiding this comment. Choose a reason for hiding this commentThe 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() |
There was a problem hiding this comment.
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?