root / ETSN / MySteps_2.py @ 271
Historique | Voir | Annoter | Télécharger (5,71 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 | 268 | equemene | return(res_np)
|
78 | 268 | equemene | |
79 | 268 | equemene | # OpenCL complete operation
|
80 | 268 | equemene | def OpenCLSillyAddition(a_np,b_np): |
81 | 268 | equemene | |
82 | 268 | equemene | # Context creation
|
83 | 268 | equemene | ctx = cl.create_some_context() |
84 | 268 | equemene | # Every process is stored in a queue
|
85 | 268 | equemene | queue = cl.CommandQueue(ctx) |
86 | 268 | equemene | |
87 | 268 | equemene | TimeIn=time.time() |
88 | 268 | equemene | # Copy from Host to Device using pointers
|
89 | 268 | equemene | mf = cl.mem_flags |
90 | 268 | equemene | a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np) |
91 | 268 | equemene | b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np) |
92 | 268 | equemene | Elapsed=time.time()-TimeIn |
93 | 268 | equemene | print("Copy from Host 2 Device : %.3f" % Elapsed)
|
94 | 268 | equemene | |
95 | 268 | equemene | TimeIn=time.time() |
96 | 268 | equemene | # Definition of kernel under OpenCL
|
97 | 268 | equemene | prg = cl.Program(ctx, """
|
98 | 268 | equemene |
|
99 | 268 | equemene | float MySillyFunction(float x)
|
100 | 268 | equemene | {
|
101 | 268 | equemene | return(pow(sqrt(log(exp(atanh(tanh(asinh(sinh(acosh(cosh(atan(tan(asin(sin(acos(cos(x))))))))))))))),2));
|
102 | 268 | equemene | }
|
103 | 268 | equemene |
|
104 | 268 | equemene | __kernel void sillysum(
|
105 | 268 | equemene | __global const float *a_g, __global const float *b_g, __global float *res_g)
|
106 | 268 | equemene | {
|
107 | 268 | equemene | int gid = get_global_id(0);
|
108 | 268 | equemene | res_g[gid] = MySillyFunction(a_g[gid]) + MySillyFunction(b_g[gid]);
|
109 | 268 | equemene | }
|
110 | 268 | equemene |
|
111 | 268 | equemene | __kernel void sum(
|
112 | 268 | equemene | __global const float *a_g, __global const float *b_g, __global float *res_g)
|
113 | 268 | equemene | {
|
114 | 268 | equemene | int gid = get_global_id(0);
|
115 | 268 | equemene | res_g[gid] = a_g[gid] + b_g[gid];
|
116 | 268 | equemene | }
|
117 | 268 | equemene | """).build()
|
118 | 268 | equemene | Elapsed=time.time()-TimeIn |
119 | 268 | equemene | print("Building kernels : %.3f" % Elapsed)
|
120 | 268 | equemene | |
121 | 268 | equemene | TimeIn=time.time() |
122 | 268 | equemene | # Memory allocation on Device for result
|
123 | 268 | equemene | res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes) |
124 | 268 | equemene | Elapsed=time.time()-TimeIn |
125 | 268 | equemene | print("Allocation on Device for results : %.3f" % Elapsed)
|
126 | 268 | equemene | |
127 | 268 | equemene | TimeIn=time.time() |
128 | 268 | equemene | # Synthesis of function "sillysum" inside Kernel Sources
|
129 | 268 | equemene | knl = prg.sillysum # Use this Kernel object for repeated calls
|
130 | 268 | equemene | Elapsed=time.time()-TimeIn |
131 | 268 | equemene | print("Synthesis of kernel : %.3f" % Elapsed)
|
132 | 268 | equemene | |
133 | 268 | equemene | TimeIn=time.time() |
134 | 268 | equemene | # Call of kernel previously defined
|
135 | 268 | equemene | CallCL=knl(queue, a_np.shape, None, a_g, b_g, res_g)
|
136 | 268 | equemene | #
|
137 | 268 | equemene | CallCL.wait() |
138 | 268 | equemene | Elapsed=time.time()-TimeIn |
139 | 268 | equemene | print("Execution of kernel : %.3f" % Elapsed)
|
140 | 268 | equemene | |
141 | 268 | equemene | TimeIn=time.time() |
142 | 268 | equemene | # Creation of vector for result with same size as input vectors
|
143 | 268 | equemene | res_np = np.empty_like(a_np) |
144 | 268 | equemene | Elapsed=time.time()-TimeIn |
145 | 268 | equemene | print("Allocation on Host for results: %.3f" % Elapsed)
|
146 | 268 | equemene | |
147 | 268 | equemene | TimeIn=time.time() |
148 | 268 | equemene | # Copy from Device to Host
|
149 | 268 | equemene | cl.enqueue_copy(queue, res_np, res_g) |
150 | 268 | equemene | Elapsed=time.time()-TimeIn |
151 | 268 | equemene | print("Copy from Device 2 Host : %.3f" % Elapsed)
|
152 | 268 | equemene | |
153 | 268 | equemene | return(res_np)
|
154 | 268 | equemene | |
155 | 268 | equemene | import sys |
156 | 268 | equemene | import time |
157 | 268 | equemene | |
158 | 268 | equemene | if __name__=='__main__': |
159 | 268 | equemene | |
160 | 268 | equemene | # Size of input vectors definition based on stdin
|
161 | 268 | equemene | import sys |
162 | 268 | equemene | try:
|
163 | 268 | equemene | SIZE=int(sys.argv[1]) |
164 | 268 | equemene | print("Size of vectors set to %i" % SIZE)
|
165 | 268 | equemene | except:
|
166 | 268 | equemene | SIZE=50000
|
167 | 268 | equemene | print("Size of vectors set to default size %i" % SIZE)
|
168 | 268 | equemene | |
169 | 268 | equemene | a_np = np.random.rand(SIZE).astype(np.float32) |
170 | 268 | equemene | b_np = np.random.rand(SIZE).astype(np.float32) |
171 | 268 | equemene | |
172 | 268 | equemene | TimeIn=time.time() |
173 | 268 | equemene | res_np=NativeSillyAddition(a_np,b_np) |
174 | 268 | equemene | NativeElapsed=time.time()-TimeIn |
175 | 268 | equemene | NativeRate=int(SIZE/NativeElapsed)
|
176 | 268 | equemene | print("NativeRate: %i" % NativeRate)
|
177 | 268 | equemene | |
178 | 268 | equemene | TimeIn=time.time() |
179 | 268 | equemene | res_cl=OpenCLSillyAddition(a_np,b_np) |
180 | 268 | equemene | OpenCLElapsed=time.time()-TimeIn |
181 | 268 | equemene | OpenCLRate=int(SIZE/OpenCLElapsed)
|
182 | 268 | equemene | print("OpenCLRate: %i" % OpenCLRate)
|
183 | 268 | equemene | |
184 | 268 | equemene | print("OpenCLvsNative ratio: %f" % (OpenCLRate/NativeRate))
|
185 | 268 | equemene | |
186 | 268 | equemene | # Check on CPU with Numpy:
|
187 | 268 | equemene | print(res_cl - res_np) |
188 | 268 | equemene | print(np.linalg.norm(res_cl - res_np)) |
189 | 268 | equemene | assert np.allclose(res_np, res_cl,rtol=1e-4) |