Skip to content

Instantly share code, notes, and snippets.

@hughperkins
Created November 5, 2016 11:21
Show Gist options
  • Select an option

  • Save hughperkins/6d9d010225f8c9a5e28a7526c7063e78 to your computer and use it in GitHub Desktop.

Select an option

Save hughperkins/6d9d010225f8c9a5e28a7526c7063e78 to your computer and use it in GitHub Desktop.
import numpy as np
import pyopencl as cl
N = 32
its = 3
np.random.seed(123)
a = np.random.rand(N).astype(np.float32)
gpu_idx = 0
platforms = cl.get_platforms()
i = 0
for platform in platforms:
gpu_devices = platform.get_devices(device_type=cl.device_type.GPU)
if gpu_idx < i + len(gpu_devices):
ctx = cl.Context(devices=[gpu_devices[gpu_idx - i]])
break
i += len(gpu_devices)
print('context', ctx)
q = cl.CommandQueue(ctx)
mf = cl.mem_flags
a_gpu = cl.Buffer(ctx, mf.READ_WRITE | mf.COPY_HOST_PTR, hostbuf=a)
prg = cl.Program(ctx, """
kernel void _z8setValuePfif(global float* data, long data_offset, int idx, float value) {
data = (global float*)((global char *)data + data_offset);
label0:;
int v1 = get_local_id(0);
bool v2 = v1 == 0;
if(v2) {
goto v4;
} else {
goto v5;
}
v4:;
long v6 = idx;
global float* v7 = (&data[v6]);
v7[0] = value;
goto v5;
v5:;
return;
}
""").build()
print('run kernel...')
workgroupsize = 32
global_size = ((N + workgroupsize - 1) // workgroupsize) * workgroupsize
cl.enqueue_copy(q, a_gpu, a)
for it in range(its):
prg._z8setValuePfif(q, (global_size,), (workgroupsize,), a_gpu, np.int64(0), np.int32(0), np.float32(123))
a_res = np.empty_like(a)
cl.enqueue_copy(q, a_res, a_gpu)
q.finish()
print('kernel done')
print('a_res[:5]', a_res[:5])
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment