1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
|
import pyopencl as cl
mf = cl.mem_flags
import numpy
class Memory:
def __init__(self, descriptor, geometry, context, float_type):
self.context = context
self.float_type = float_type
self.size_x = geometry.size_x
self.size_y = geometry.size_y
self.size_z = geometry.size_z
self.volume = self.size_x * self.size_y * self.size_z
self.pop_size = descriptor.q * self.volume * self.float_type(0).nbytes
self.moments_size = 3 * self.volume * self.float_type(0).nbytes
self.cl_pop = cl.Buffer(self.context, mf.READ_WRITE, size=self.pop_size)
self.cl_moments = cl.Buffer(self.context, mf.WRITE_ONLY, size=self.moments_size)
def gid(self, x, y, z = 0):
return z * (self.size_x*self.size_y) + y * self.size_x + x;
class Lattice:
def __init__(self, geometry, kernel_src, descriptor, platform = 0, precision = 'single'):
self.geometry = geometry
self.descriptor = descriptor
self.float_type = {
'single': (numpy.float32, 'float'),
'double': (numpy.float64, 'double'),
}.get(precision, None)
self.platform = cl.get_platforms()[platform]
self.layout = None
self.context = cl.Context(
properties=[(cl.context_properties.PLATFORM, self.platform)])
self.queue = cl.CommandQueue(self.context)
self.memory = Memory(descriptor, self.geometry, self.context, self.float_type[0])
self.tick = False
self.compiler_args = {
'single': '-cl-single-precision-constant -cl-fast-relaxed-math',
'double': '-cl-fast-relaxed-math'
}.get(precision, None)
self.build_kernel(kernel_src)
self.program.equilibrilize_all(
self.queue, self.geometry.size(), self.layout, self.memory.cl_pop).wait()
self.tick_tasks = []
self.tock_tasks = []
def build_kernel(self, src):
self.program = cl.Program(self.context, src).build(self.compiler_args)
def schedule_tick(self, f, cells, *params):
self.tick_tasks += [ (eval("self.program.%s" % f), cells, params) ]
def schedule_tock(self, f, cells, *params):
self.tock_tasks += [ (eval("self.program.%s" % f), cells, params) ]
def evolve(self):
if self.tick:
self.tick = False
for f, cells, params in self.tick_tasks:
f(self.queue, cells.size(), self.layout, self.memory.cl_pop, cells.get(), *params)
else:
self.tick = True
for f, cells, params in self.tock_tasks:
f(self.queue, cells.size(), self.layout, self.memory.cl_pop, cells.get(), *params)
def sync(self):
self.queue.finish()
def get_moments(self):
moments = numpy.ndarray(shape=(self.memory.volume*(self.descriptor.d+1),1), dtype=self.float_type[0])
self.program.collect_moments_all(
self.queue, self.geometry.inner_size(), self.layout, self.memory.cl_pop, self.memory.cl_moments)
cl.enqueue_copy(self.queue, moments, self.memory.cl_moments).wait();
return moments
HelperTemplate = """
__kernel void equilibrilize_all(__global ${float_type}* f_next)
{
const unsigned int gid = ${index.gid('get_global_id(0)', 'get_global_id(1)')};
equilibrilize_tick(f_next, gid);
}
__kernel void collect_moments_all(__global ${float_type}* f,
__global ${float_type}* moments)
{
const unsigned int gid = ${index.gid('get_global_id(0)+1', 'get_global_id(1)+1')};
collect_moments_tock(f, gid, moments);
}
"""
|
