root / ETSN / MySteps_0.py @ 272
Historique | Voir | Annoter | Télécharger (1,62 ko)
1 |
#!/usr/bin/env python3
|
---|---|
2 |
|
3 |
import numpy as np |
4 |
import pyopencl as cl |
5 |
|
6 |
# Native Operation under Numpy (for prototyping & tests
|
7 |
def NativeAddition(a_np,b_np): |
8 |
return(a_np+b_np)
|
9 |
|
10 |
# OpenCL complete operation
|
11 |
def OpenCLAddition(a_np,b_np): |
12 |
|
13 |
# Context creation
|
14 |
ctx = cl.create_some_context() |
15 |
# Every process is stored in a queue
|
16 |
queue = cl.CommandQueue(ctx) |
17 |
|
18 |
# Copy from Host to Device using pointers
|
19 |
mf = cl.mem_flags |
20 |
a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np) |
21 |
b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np) |
22 |
|
23 |
# Definition of kernel under OpenCL
|
24 |
prg = cl.Program(ctx, """
|
25 |
__kernel void sum(
|
26 |
__global const float *a_g, __global const float *b_g, __global float *res_g)
|
27 |
{
|
28 |
int gid = get_global_id(0);
|
29 |
res_g[gid] = a_g[gid] + b_g[gid];
|
30 |
}
|
31 |
""").build()
|
32 |
|
33 |
# Memory allocation on Device for result
|
34 |
res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes) |
35 |
# Synthesis of function "sum" inside Kernel Sources
|
36 |
knl = prg.sum # Use this Kernel object for repeated calls
|
37 |
# Call of kernel previously defined
|
38 |
knl(queue, a_np.shape, None, a_g, b_g, res_g)
|
39 |
|
40 |
# Creation of vector for result with same size as input vectors
|
41 |
res_np = np.empty_like(a_np) |
42 |
# Copy from Device to Host
|
43 |
cl.enqueue_copy(queue, res_np, res_g) |
44 |
|
45 |
return(res_np)
|
46 |
|
47 |
#if __name__=='__main__':
|
48 |
|
49 |
a_np = np.random.rand(50000).astype(np.float32)
|
50 |
b_np = np.random.rand(50000).astype(np.float32)
|
51 |
|
52 |
res_np=NativeAddition(a_np,b_np) |
53 |
res_cl=OpenCLAddition(a_np,b_np) |
54 |
|
55 |
# Check on CPU with Numpy:
|
56 |
print(res_cl - res_np) |
57 |
print(np.linalg.norm(res_cl - res_np)) |
58 |
assert np.allclose(res_np, res_cl)
|