root / ETSN / MySteps_2.py @ 308
Historique | Voir | Annoter | Télécharger (5,87 ko)
1 | 268 | equemene | #!/usr/bin/env python3
|
---|---|---|---|
2 | 268 | equemene | |
3 | 268 | equemene | import numpy as np |
4 | 268 | equemene | import pyopencl as cl |
5 | 268 | equemene | |
6 | 268 | equemene | # piling 16 arithmetical functions
|
7 | 268 | equemene | def MySillyFunction(x): |
8 | 268 | equemene | return(np.power(np.sqrt(np.log(np.exp(np.arctanh(np.tanh(np.arcsinh(np.sinh(np.arccosh(np.cosh(np.arctan(np.tan(np.arcsin(np.sin(np.arccos(np.cos(x))))))))))))))),2)) |
9 | 268 | equemene | |
10 | 268 | equemene | # Native Operation under Numpy (for prototyping & tests
|
11 | 268 | equemene | def NativeAddition(a_np,b_np): |
12 | 268 | equemene | return(a_np+b_np)
|
13 | 268 | equemene | |
14 | 268 | equemene | # Native Operation with MySillyFunction under Numpy (for prototyping & tests
|
15 | 268 | equemene | def NativeSillyAddition(a_np,b_np): |
16 | 268 | equemene | return(MySillyFunction(a_np)+MySillyFunction(b_np))
|
17 | 268 | equemene | |
18 | 268 | equemene | # OpenCL complete operation
|
19 | 268 | equemene | def OpenCLAddition(a_np,b_np): |
20 | 268 | equemene | |
21 | 268 | equemene | # Context creation
|
22 | 268 | equemene | ctx = cl.create_some_context() |
23 | 268 | equemene | # Every process is stored in a queue
|
24 | 268 | equemene | queue = cl.CommandQueue(ctx) |
25 | 268 | equemene | |
26 | 268 | equemene | TimeIn=time.time() |
27 | 268 | equemene | # Copy from Host to Device using pointers
|
28 | 268 | equemene | mf = cl.mem_flags |
29 | 268 | equemene | a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np) |
30 | 268 | equemene | b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np) |
31 | 268 | equemene | Elapsed=time.time()-TimeIn |
32 | 268 | equemene | print("Copy from Host 2 Device : %.3f" % Elapsed)
|
33 | 268 | equemene | |
34 | 268 | equemene | TimeIn=time.time() |
35 | 268 | equemene | # Definition of kernel under OpenCL
|
36 | 268 | equemene | prg = cl.Program(ctx, """
|
37 | 268 | equemene | __kernel void sum(
|
38 | 268 | equemene | __global const float *a_g, __global const float *b_g, __global float *res_g)
|
39 | 268 | equemene | {
|
40 | 268 | equemene | int gid = get_global_id(0);
|
41 | 268 | equemene | res_g[gid] = a_g[gid] + b_g[gid];
|
42 | 268 | equemene | }
|
43 | 268 | equemene | """).build()
|
44 | 268 | equemene | Elapsed=time.time()-TimeIn |
45 | 268 | equemene | print("Building kernels : %.3f" % Elapsed)
|
46 | 268 | equemene | |
47 | 268 | equemene | TimeIn=time.time() |
48 | 268 | equemene | # Memory allocation on Device for result
|
49 | 268 | equemene | res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes) |
50 | 268 | equemene | Elapsed=time.time()-TimeIn |
51 | 268 | equemene | print("Allocation on Device for results : %.3f" % Elapsed)
|
52 | 268 | equemene | |
53 | 268 | equemene | TimeIn=time.time() |
54 | 268 | equemene | # Synthesis of function "sum" inside Kernel Sources
|
55 | 268 | equemene | knl = prg.sum # Use this Kernel object for repeated calls
|
56 | 268 | equemene | Elapsed=time.time()-TimeIn |
57 | 268 | equemene | print("Synthesis of kernel : %.3f" % Elapsed)
|
58 | 268 | equemene | |
59 | 268 | equemene | TimeIn=time.time() |
60 | 268 | equemene | # Call of kernel previously defined
|
61 | 268 | equemene | knl(queue, a_np.shape, None, a_g, b_g, res_g)
|
62 | 268 | equemene | Elapsed=time.time()-TimeIn |
63 | 268 | equemene | print("Execution of kernel : %.3f" % Elapsed)
|
64 | 268 | equemene | |
65 | 268 | equemene | TimeIn=time.time() |
66 | 268 | equemene | # Creation of vector for result with same size as input vectors
|
67 | 268 | equemene | res_np = np.empty_like(a_np) |
68 | 268 | equemene | Elapsed=time.time()-TimeIn |
69 | 268 | equemene | print("Allocation on Host for results: %.3f" % Elapsed)
|
70 | 268 | equemene | |
71 | 268 | equemene | TimeIn=time.time() |
72 | 268 | equemene | # Copy from Device to Host
|
73 | 268 | equemene | cl.enqueue_copy(queue, res_np, res_g) |
74 | 268 | equemene | Elapsed=time.time()-TimeIn |
75 | 268 | equemene | print("Copy from Device 2 Host : %.3f" % Elapsed)
|
76 | 268 | equemene | |
77 | 275 | equemene | # Liberation of memory
|
78 | 275 | equemene | a_g.release() |
79 | 275 | equemene | b_g.release() |
80 | 275 | equemene | res_g.release() |
81 | 275 | equemene | |
82 | 268 | equemene | return(res_np)
|
83 | 268 | equemene | |
84 | 268 | equemene | # OpenCL complete operation
|
85 | 268 | equemene | def OpenCLSillyAddition(a_np,b_np): |
86 | 268 | equemene | |
87 | 268 | equemene | # Context creation
|
88 | 268 | equemene | ctx = cl.create_some_context() |
89 | 268 | equemene | # Every process is stored in a queue
|
90 | 268 | equemene | queue = cl.CommandQueue(ctx) |
91 | 268 | equemene | |
92 | 268 | equemene | TimeIn=time.time() |
93 | 268 | equemene | # Copy from Host to Device using pointers
|
94 | 268 | equemene | mf = cl.mem_flags |
95 | 268 | equemene | a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np) |
96 | 268 | equemene | b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np) |
97 | 268 | equemene | Elapsed=time.time()-TimeIn |
98 | 268 | equemene | print("Copy from Host 2 Device : %.3f" % Elapsed)
|
99 | 268 | equemene | |
100 | 268 | equemene | TimeIn=time.time() |
101 | 268 | equemene | # Definition of kernel under OpenCL
|
102 | 268 | equemene | prg = cl.Program(ctx, """
|
103 | 268 | equemene |
|
104 | 268 | equemene | float MySillyFunction(float x)
|
105 | 268 | equemene | {
|
106 | 268 | equemene | return(pow(sqrt(log(exp(atanh(tanh(asinh(sinh(acosh(cosh(atan(tan(asin(sin(acos(cos(x))))))))))))))),2));
|
107 | 268 | equemene | }
|
108 | 268 | equemene |
|
109 | 268 | equemene | __kernel void sillysum(
|
110 | 268 | equemene | __global const float *a_g, __global const float *b_g, __global float *res_g)
|
111 | 268 | equemene | {
|
112 | 268 | equemene | int gid = get_global_id(0);
|
113 | 268 | equemene | res_g[gid] = MySillyFunction(a_g[gid]) + MySillyFunction(b_g[gid]);
|
114 | 268 | equemene | }
|
115 | 268 | equemene |
|
116 | 268 | equemene | __kernel void sum(
|
117 | 268 | equemene | __global const float *a_g, __global const float *b_g, __global float *res_g)
|
118 | 268 | equemene | {
|
119 | 268 | equemene | int gid = get_global_id(0);
|
120 | 268 | equemene | res_g[gid] = a_g[gid] + b_g[gid];
|
121 | 268 | equemene | }
|
122 | 268 | equemene | """).build()
|
123 | 268 | equemene | Elapsed=time.time()-TimeIn |
124 | 268 | equemene | print("Building kernels : %.3f" % Elapsed)
|
125 | 268 | equemene | |
126 | 268 | equemene | TimeIn=time.time() |
127 | 268 | equemene | # Memory allocation on Device for result
|
128 | 268 | equemene | res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes) |
129 | 268 | equemene | Elapsed=time.time()-TimeIn |
130 | 268 | equemene | print("Allocation on Device for results : %.3f" % Elapsed)
|
131 | 268 | equemene | |
132 | 268 | equemene | TimeIn=time.time() |
133 | 268 | equemene | # Synthesis of function "sillysum" inside Kernel Sources
|
134 | 268 | equemene | knl = prg.sillysum # Use this Kernel object for repeated calls
|
135 | 268 | equemene | Elapsed=time.time()-TimeIn |
136 | 268 | equemene | print("Synthesis of kernel : %.3f" % Elapsed)
|
137 | 268 | equemene | |
138 | 268 | equemene | TimeIn=time.time() |
139 | 268 | equemene | # Call of kernel previously defined
|
140 | 268 | equemene | CallCL=knl(queue, a_np.shape, None, a_g, b_g, res_g)
|
141 | 268 | equemene | #
|
142 | 268 | equemene | CallCL.wait() |
143 | 268 | equemene | Elapsed=time.time()-TimeIn |
144 | 268 | equemene | print("Execution of kernel : %.3f" % Elapsed)
|
145 | 268 | equemene | |
146 | 268 | equemene | TimeIn=time.time() |
147 | 268 | equemene | # Creation of vector for result with same size as input vectors
|
148 | 268 | equemene | res_np = np.empty_like(a_np) |
149 | 268 | equemene | Elapsed=time.time()-TimeIn |
150 | 268 | equemene | print("Allocation on Host for results: %.3f" % Elapsed)
|
151 | 268 | equemene | |
152 | 268 | equemene | TimeIn=time.time() |
153 | 268 | equemene | # Copy from Device to Host
|
154 | 268 | equemene | cl.enqueue_copy(queue, res_np, res_g) |
155 | 268 | equemene | Elapsed=time.time()-TimeIn |
156 | 268 | equemene | print("Copy from Device 2 Host : %.3f" % Elapsed)
|
157 | 268 | equemene | |
158 | 275 | equemene | # Liberation of memory
|
159 | 275 | equemene | a_g.release() |
160 | 275 | equemene | b_g.release() |
161 | 275 | equemene | res_g.release() |
162 | 275 | equemene | |
163 | 268 | equemene | return(res_np)
|
164 | 268 | equemene | |
165 | 268 | equemene | import sys |
166 | 268 | equemene | import time |
167 | 268 | equemene | |
168 | 268 | equemene | if __name__=='__main__': |
169 | 268 | equemene | |
170 | 268 | equemene | # Size of input vectors definition based on stdin
|
171 | 268 | equemene | import sys |
172 | 268 | equemene | try:
|
173 | 268 | equemene | SIZE=int(sys.argv[1]) |
174 | 268 | equemene | print("Size of vectors set to %i" % SIZE)
|
175 | 268 | equemene | except:
|
176 | 268 | equemene | SIZE=50000
|
177 | 268 | equemene | print("Size of vectors set to default size %i" % SIZE)
|
178 | 268 | equemene | |
179 | 268 | equemene | a_np = np.random.rand(SIZE).astype(np.float32) |
180 | 268 | equemene | b_np = np.random.rand(SIZE).astype(np.float32) |
181 | 268 | equemene | |
182 | 268 | equemene | TimeIn=time.time() |
183 | 268 | equemene | res_np=NativeSillyAddition(a_np,b_np) |
184 | 268 | equemene | NativeElapsed=time.time()-TimeIn |
185 | 268 | equemene | NativeRate=int(SIZE/NativeElapsed)
|
186 | 268 | equemene | print("NativeRate: %i" % NativeRate)
|
187 | 268 | equemene | |
188 | 268 | equemene | TimeIn=time.time() |
189 | 268 | equemene | res_cl=OpenCLSillyAddition(a_np,b_np) |
190 | 268 | equemene | OpenCLElapsed=time.time()-TimeIn |
191 | 268 | equemene | OpenCLRate=int(SIZE/OpenCLElapsed)
|
192 | 268 | equemene | print("OpenCLRate: %i" % OpenCLRate)
|
193 | 268 | equemene | |
194 | 268 | equemene | print("OpenCLvsNative ratio: %f" % (OpenCLRate/NativeRate))
|
195 | 268 | equemene | |
196 | 268 | equemene | # Check on CPU with Numpy:
|
197 | 268 | equemene | print(res_cl - res_np) |
198 | 268 | equemene | print(np.linalg.norm(res_cl - res_np)) |
199 | 296 | equemene | assert np.allclose(res_cl, res_np,rtol=1e-4) |