Révision 270

ETSN/MyDFT_1.py (revision 270)
1
#!/usr/bin/env python3
2

  
3
import numpy as np
4
import pyopencl as cl
5

  
6
# piling 16 arithmetical functions
7
def MySillyFunction(x):
8
    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

  
10
# Native Operation under Numpy (for prototyping & tests
11
def NativeAddition(a_np,b_np):
12
    return(a_np+b_np)
13

  
14
# Native Operation with MySillyFunction under Numpy (for prototyping & tests
15
def NativeSillyAddition(a_np,b_np):
16
    return(MySillyFunction(a_np)+MySillyFunction(b_np))
17

  
18
# Naive Discrete Fourier Transform
19
def MyDFT(x,y):
20
    from numpy import pi,cos,sin
21
    size=x.shape[0]
22
    X=np.zeros(size).astype(np.float32)
23
    Y=np.zeros(size).astype(np.float32)
24
    for i in range(size):
25
        for j in range(size):
26
            X[i]=X[i]+x[j]*cos(2.*pi*i*j/size)-y[j]*sin(2.*pi*i*j/size)
27
            Y[i]=Y[i]+x[j]*sin(2.*pi*i*j/size)+y[j]*cos(2.*pi*i*j/size)
28
    return(X,Y)
29

  
30
# CUDA complete operation
31
def CUDAAddition(a_np,b_np):
32
    import pycuda.autoinit
33
    import pycuda.driver as drv
34
    import numpy
35

  
36
    from pycuda.compiler import SourceModule
37
    mod = SourceModule("""
38
    __global__ void sum(float *dest, float *a, float *b)
39
{
40
  // const int i = threadIdx.x;
41
  const int i = blockIdx.x;
42
  dest[i] = a[i] + b[i];
43
}
44
""")
45

  
46
    # sum = mod.get_function("sum")
47
    sum = mod.get_function("sum")
48

  
49
    res_np = numpy.zeros_like(a_np)
50
    sum(drv.Out(res_np), drv.In(a_np), drv.In(b_np),
51
        block=(1,1,1), grid=(a_np.size,1))
52
    return(res_np)
53

  
54
# CUDA Silly complete operation
55
def CUDASillyAddition(a_np,b_np):
56
    import pycuda.autoinit
57
    import pycuda.driver as drv
58
    import numpy
59

  
60
    from pycuda.compiler import SourceModule
61
    TimeIn=time.time()
62
    mod = SourceModule("""
63
__device__ float MySillyFunction(float x)
64
{
65
    return(pow(sqrt(log(exp(atanh(tanh(asinh(sinh(acosh(cosh(atan(tan(asin(sin(acos(cos(x))))))))))))))),2)); 
66
}
67

  
68
__global__ void sillysum(float *dest, float *a, float *b)
69
{
70
  const int i = blockIdx.x;
71
  dest[i] = MySillyFunction(a[i]) + MySillyFunction(b[i]);
72
}
73
""")
74
    Elapsed=time.time()-TimeIn
75
    print("Definition of kernel : %.3f" % Elapsed)
76

  
77
    TimeIn=time.time()
78
    # sum = mod.get_function("sum")
79
    sillysum = mod.get_function("sillysum")
80
    Elapsed=time.time()-TimeIn
81
    print("Synthesis of kernel : %.3f" % Elapsed)
82

  
83
    TimeIn=time.time()
84
    res_np = numpy.zeros_like(a_np)
85
    Elapsed=time.time()-TimeIn
86
    print("Allocation on Host for results : %.3f" % Elapsed)
87

  
88
    TimeIn=time.time()
89
    sillysum(drv.Out(res_np), drv.In(a_np), drv.In(b_np),
90
             block=(1,1,1), grid=(a_np.size,1))
91
    Elapsed=time.time()-TimeIn
92
    print("Execution of kernel : %.3f" % Elapsed)
93
    return(res_np)
94

  
95
# OpenCL complete operation
96
def OpenCLAddition(a_np,b_np):
97

  
98
    # Context creation
99
    ctx = cl.create_some_context()
100
    # Every process is stored in a queue
101
    queue = cl.CommandQueue(ctx)
102

  
103
    TimeIn=time.time()
104
    # Copy from Host to Device using pointers
105
    mf = cl.mem_flags
106
    a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np)
107
    b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np)
108
    Elapsed=time.time()-TimeIn
109
    print("Copy from Host 2 Device : %.3f" % Elapsed)
110

  
111
    TimeIn=time.time()
112
    # Definition of kernel under OpenCL
113
    prg = cl.Program(ctx, """
114
__kernel void sum(
115
    __global const float *a_g, __global const float *b_g, __global float *res_g)
116
{
117
  int gid = get_global_id(0);
118
  res_g[gid] = a_g[gid] + b_g[gid];
119
}
120
""").build()
121
    Elapsed=time.time()-TimeIn
122
    print("Building kernels : %.3f" % Elapsed)
123
    
124
    TimeIn=time.time()
125
    # Memory allocation on Device for result
126
    res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes)
127
    Elapsed=time.time()-TimeIn
128
    print("Allocation on Device for results : %.3f" % Elapsed)
129

  
130
    TimeIn=time.time()
131
    # Synthesis of function "sum" inside Kernel Sources
132
    knl = prg.sum  # Use this Kernel object for repeated calls
133
    Elapsed=time.time()-TimeIn
134
    print("Synthesis of kernel : %.3f" % Elapsed)
135

  
136
    TimeIn=time.time()
137
    # Call of kernel previously defined 
138
    knl(queue, a_np.shape, None, a_g, b_g, res_g)
139
    Elapsed=time.time()-TimeIn
140
    print("Execution of kernel : %.3f" % Elapsed)
141

  
142
    TimeIn=time.time()
143
    # Creation of vector for result with same size as input vectors
144
    res_np = np.empty_like(a_np)
145
    Elapsed=time.time()-TimeIn
146
    print("Allocation on Host for results: %.3f" % Elapsed)
147

  
148
    TimeIn=time.time()
149
    # Copy from Device to Host
150
    cl.enqueue_copy(queue, res_np, res_g)
151
    Elapsed=time.time()-TimeIn
152
    print("Copy from Device 2 Host : %.3f" % Elapsed)
153

  
154
    return(res_np)
155

  
156
# OpenCL complete operation
157
def OpenCLSillyAddition(a_np,b_np):
158

  
159
    # Context creation
160
    ctx = cl.create_some_context()
161
    # Every process is stored in a queue
162
    queue = cl.CommandQueue(ctx)
163

  
164
    TimeIn=time.time()
165
    # Copy from Host to Device using pointers
166
    mf = cl.mem_flags
167
    a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np)
168
    b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np)
169
    Elapsed=time.time()-TimeIn
170
    print("Copy from Host 2 Device : %.3f" % Elapsed)
171

  
172
    TimeIn=time.time()
173
    # Definition of kernel under OpenCL
174
    prg = cl.Program(ctx, """
175

  
176
float MySillyFunction(float x)
177
{
178
    return(pow(sqrt(log(exp(atanh(tanh(asinh(sinh(acosh(cosh(atan(tan(asin(sin(acos(cos(x))))))))))))))),2)); 
179
}
180

  
181
__kernel void sillysum(
182
    __global const float *a_g, __global const float *b_g, __global float *res_g)
183
{
184
  int gid = get_global_id(0);
185
  res_g[gid] = MySillyFunction(a_g[gid]) + MySillyFunction(b_g[gid]);
186
}
187

  
188
__kernel void sum(
189
    __global const float *a_g, __global const float *b_g, __global float *res_g)
190
{
191
  int gid = get_global_id(0);
192
  res_g[gid] = a_g[gid] + b_g[gid];
193
}
194
""").build()
195
    Elapsed=time.time()-TimeIn
196
    print("Building kernels : %.3f" % Elapsed)
197
    
198
    TimeIn=time.time()
199
    # Memory allocation on Device for result
200
    res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes)
201
    Elapsed=time.time()-TimeIn
202
    print("Allocation on Device for results : %.3f" % Elapsed)
203

  
204
    TimeIn=time.time()
205
    # Synthesis of function "sillysum" inside Kernel Sources
206
    knl = prg.sillysum  # Use this Kernel object for repeated calls
207
    Elapsed=time.time()-TimeIn
208
    print("Synthesis of kernel : %.3f" % Elapsed)
209

  
210
    TimeIn=time.time()
211
    # Call of kernel previously defined 
212
    CallCL=knl(queue, a_np.shape, None, a_g, b_g, res_g)
213
    # 
214
    CallCL.wait()
215
    Elapsed=time.time()-TimeIn
216
    print("Execution of kernel : %.3f" % Elapsed)
217

  
218
    TimeIn=time.time()
219
    # Creation of vector for result with same size as input vectors
220
    res_np = np.empty_like(a_np)
221
    Elapsed=time.time()-TimeIn
222
    print("Allocation on Host for results: %.3f" % Elapsed)
223

  
224
    TimeIn=time.time()
225
    # Copy from Device to Host
226
    cl.enqueue_copy(queue, res_np, res_g)
227
    Elapsed=time.time()-TimeIn
228
    print("Copy from Device 2 Host : %.3f" % Elapsed)
229

  
230
    return(res_np)
231

  
232
import sys
233
import time
234

  
235
if __name__=='__main__':
236

  
237
    # Size of input vectors definition based on stdin
238
    import sys
239
    try:
240
        SIZE=int(sys.argv[1])
241
        print("Size of vectors set to %i" % SIZE)
242
    except: 
243
        SIZE=50000
244
        print("Size of vectors set to default size %i" % SIZE)
245
        
246
    # a_np = np.random.rand(SIZE).astype(np.float32)
247
    # b_np = np.random.rand(SIZE).astype(np.float32)
248
    
249
    a_np = np.ones(SIZE).astype(np.float32)
250
    b_np = np.ones(SIZE).astype(np.float32)
251

  
252
    # Native Implementation
253
    TimeIn=time.time()
254
    # res_np=NativeAddition(a_np,b_np)
255
    # res_np=NativeSillyAddition(a_np,b_np)
256
    c_np,d_np=MyDFT(a_np,b_np)
257
    NativeElapsed=time.time()-TimeIn
258
    NativeRate=int(SIZE/NativeElapsed)
259
    print("NativeRate: %i" % NativeRate)
260

  
261
    print(c_np,d_np) 
262
    
263
   #  # OpenCL Implementation
264
   #  TimeIn=time.time()
265
   #  # res_cl=OpenCLAddition(a_np,b_np)
266
   #  res_cl=OpenCLSillyAddition(a_np,b_np)
267
   #  OpenCLElapsed=time.time()-TimeIn
268
   #  OpenCLRate=int(SIZE/OpenCLElapsed)
269
   #  print("OpenCLRate: %i" % OpenCLRate)
270

  
271
   #  # CUDA Implementation
272
   #  TimeIn=time.time()
273
   #  # res_cuda=CUDAAddition(a_np,b_np)
274
   #  res_cuda=CUDASillyAddition(a_np,b_np)
275
   #  CUDAElapsed=time.time()-TimeIn
276
   #  CUDARate=int(SIZE/CUDAElapsed)
277
   #  print("CUDARate: %i" % CUDARate)
278
    
279
   #  print("OpenCLvsNative ratio: %f" % (OpenCLRate/NativeRate))
280
   #  print("CUDAvsNative ratio: %f" % (CUDARate/NativeRate))
281
    
282
   # # Check on OpenCL with Numpy:
283
   #  print(res_cl - res_np)
284
   #  print(np.linalg.norm(res_cl - res_np))
285
   #  try:
286
   #      assert np.allclose(res_np, res_cl)
287
   #  except:
288
   #      print("Results between Native & OpenCL seem to be too different!")
289
        
290
   #  # Check on CUDA with Numpy:
291
   #  print(res_cuda - res_np)
292
   #  print(np.linalg.norm(res_cuda - res_np))
293
   #  try:
294
   #      assert np.allclose(res_np, res_cuda)
295
   #  except:
296
   #      print("Results between Native & CUDA seem to be too different!")
297

  
298

  
0 299

  
ETSN/MyDFT_2.py (revision 270)
1
#!/usr/bin/env python3
2

  
3
import numpy as np
4
import pyopencl as cl
5

  
6
# piling 16 arithmetical functions
7
def MySillyFunction(x):
8
    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

  
10
# Native Operation under Numpy (for prototyping & tests
11
def NativeAddition(a_np,b_np):
12
    return(a_np+b_np)
13

  
14
# Native Operation with MySillyFunction under Numpy (for prototyping & tests
15
def NativeSillyAddition(a_np,b_np):
16
    return(MySillyFunction(a_np)+MySillyFunction(b_np))
17

  
18
# Naive Discrete Fourier Transform
19
def MyDFT(x,y):
20
    from numpy import pi,cos,sin
21
    size=x.shape[0]
22
    X=np.zeros(size).astype(np.float32)
23
    Y=np.zeros(size).astype(np.float32)
24
    for i in range(size):
25
        for j in range(size):
26
            X[i]=X[i]+x[j]*cos(2.*pi*i*j/size)-y[j]*sin(2.*pi*i*j/size)
27
            Y[i]=Y[i]+x[j]*sin(2.*pi*i*j/size)+y[j]*cos(2.*pi*i*j/size)
28
    return(X,Y)
29

  
30
# Numpy Discrete Fourier Transform
31
def NumpyDFT(x,y):
32
    from numpy import pi,cos,sin
33
    size=x.shape[0]
34
    X=np.zeros(size).astype(np.float32)
35
    Y=np.zeros(size).astype(np.float32)
36
    nj=np.multiply(2.0*np.pi/size,np.arange(size)).astype(np.float32)
37
    for i in range(size):
38
        X[i]=np.sum(np.subtract(np.multiply(np.cos(i*nj),x),np.multiply(np.sin(i*nj),y)))
39
        Y[i]=np.sum(np.add(np.multiply(np.sin(i*nj),x),np.multiply(np.cos(i*nj),y)))
40
    return(X,Y)
41

  
42
# CUDA complete operation
43
def CUDAAddition(a_np,b_np):
44
    import pycuda.autoinit
45
    import pycuda.driver as drv
46
    import numpy
47

  
48
    from pycuda.compiler import SourceModule
49
    mod = SourceModule("""
50
    __global__ void sum(float *dest, float *a, float *b)
51
{
52
  // const int i = threadIdx.x;
53
  const int i = blockIdx.x;
54
  dest[i] = a[i] + b[i];
55
}
56
""")
57

  
58
    # sum = mod.get_function("sum")
59
    sum = mod.get_function("sum")
60

  
61
    res_np = numpy.zeros_like(a_np)
62
    sum(drv.Out(res_np), drv.In(a_np), drv.In(b_np),
63
        block=(1,1,1), grid=(a_np.size,1))
64
    return(res_np)
65

  
66
# CUDA Silly complete operation
67
def CUDASillyAddition(a_np,b_np):
68
    import pycuda.autoinit
69
    import pycuda.driver as drv
70
    import numpy
71

  
72
    from pycuda.compiler import SourceModule
73
    TimeIn=time.time()
74
    mod = SourceModule("""
75
__device__ float MySillyFunction(float x)
76
{
77
    return(pow(sqrt(log(exp(atanh(tanh(asinh(sinh(acosh(cosh(atan(tan(asin(sin(acos(cos(x))))))))))))))),2)); 
78
}
79

  
80
__global__ void sillysum(float *dest, float *a, float *b)
81
{
82
  const int i = blockIdx.x;
83
  dest[i] = MySillyFunction(a[i]) + MySillyFunction(b[i]);
84
}
85
""")
86
    Elapsed=time.time()-TimeIn
87
    print("Definition of kernel : %.3f" % Elapsed)
88

  
89
    TimeIn=time.time()
90
    # sum = mod.get_function("sum")
91
    sillysum = mod.get_function("sillysum")
92
    Elapsed=time.time()-TimeIn
93
    print("Synthesis of kernel : %.3f" % Elapsed)
94

  
95
    TimeIn=time.time()
96
    res_np = numpy.zeros_like(a_np)
97
    Elapsed=time.time()-TimeIn
98
    print("Allocation on Host for results : %.3f" % Elapsed)
99

  
100
    TimeIn=time.time()
101
    sillysum(drv.Out(res_np), drv.In(a_np), drv.In(b_np),
102
             block=(1,1,1), grid=(a_np.size,1))
103
    Elapsed=time.time()-TimeIn
104
    print("Execution of kernel : %.3f" % Elapsed)
105
    return(res_np)
106

  
107
# OpenCL complete operation
108
def OpenCLAddition(a_np,b_np):
109

  
110
    # Context creation
111
    ctx = cl.create_some_context()
112
    # Every process is stored in a queue
113
    queue = cl.CommandQueue(ctx)
114

  
115
    TimeIn=time.time()
116
    # Copy from Host to Device using pointers
117
    mf = cl.mem_flags
118
    a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np)
119
    b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np)
120
    Elapsed=time.time()-TimeIn
121
    print("Copy from Host 2 Device : %.3f" % Elapsed)
122

  
123
    TimeIn=time.time()
124
    # Definition of kernel under OpenCL
125
    prg = cl.Program(ctx, """
126
__kernel void sum(
127
    __global const float *a_g, __global const float *b_g, __global float *res_g)
128
{
129
  int gid = get_global_id(0);
130
  res_g[gid] = a_g[gid] + b_g[gid];
131
}
132
""").build()
133
    Elapsed=time.time()-TimeIn
134
    print("Building kernels : %.3f" % Elapsed)
135
    
136
    TimeIn=time.time()
137
    # Memory allocation on Device for result
138
    res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes)
139
    Elapsed=time.time()-TimeIn
140
    print("Allocation on Device for results : %.3f" % Elapsed)
141

  
142
    TimeIn=time.time()
143
    # Synthesis of function "sum" inside Kernel Sources
144
    knl = prg.sum  # Use this Kernel object for repeated calls
145
    Elapsed=time.time()-TimeIn
146
    print("Synthesis of kernel : %.3f" % Elapsed)
147

  
148
    TimeIn=time.time()
149
    # Call of kernel previously defined 
150
    knl(queue, a_np.shape, None, a_g, b_g, res_g)
151
    Elapsed=time.time()-TimeIn
152
    print("Execution of kernel : %.3f" % Elapsed)
153

  
154
    TimeIn=time.time()
155
    # Creation of vector for result with same size as input vectors
156
    res_np = np.empty_like(a_np)
157
    Elapsed=time.time()-TimeIn
158
    print("Allocation on Host for results: %.3f" % Elapsed)
159

  
160
    TimeIn=time.time()
161
    # Copy from Device to Host
162
    cl.enqueue_copy(queue, res_np, res_g)
163
    Elapsed=time.time()-TimeIn
164
    print("Copy from Device 2 Host : %.3f" % Elapsed)
165

  
166
    return(res_np)
167

  
168
# OpenCL complete operation
169
def OpenCLSillyAddition(a_np,b_np):
170

  
171
    # Context creation
172
    ctx = cl.create_some_context()
173
    # Every process is stored in a queue
174
    queue = cl.CommandQueue(ctx)
175

  
176
    TimeIn=time.time()
177
    # Copy from Host to Device using pointers
178
    mf = cl.mem_flags
179
    a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np)
180
    b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np)
181
    Elapsed=time.time()-TimeIn
182
    print("Copy from Host 2 Device : %.3f" % Elapsed)
183

  
184
    TimeIn=time.time()
185
    # Definition of kernel under OpenCL
186
    prg = cl.Program(ctx, """
187

  
188
float MySillyFunction(float x)
189
{
190
    return(pow(sqrt(log(exp(atanh(tanh(asinh(sinh(acosh(cosh(atan(tan(asin(sin(acos(cos(x))))))))))))))),2)); 
191
}
192

  
193
__kernel void sillysum(
194
    __global const float *a_g, __global const float *b_g, __global float *res_g)
195
{
196
  int gid = get_global_id(0);
197
  res_g[gid] = MySillyFunction(a_g[gid]) + MySillyFunction(b_g[gid]);
198
}
199

  
200
__kernel void sum(
201
    __global const float *a_g, __global const float *b_g, __global float *res_g)
202
{
203
  int gid = get_global_id(0);
204
  res_g[gid] = a_g[gid] + b_g[gid];
205
}
206
""").build()
207
    Elapsed=time.time()-TimeIn
208
    print("Building kernels : %.3f" % Elapsed)
209
    
210
    TimeIn=time.time()
211
    # Memory allocation on Device for result
212
    res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes)
213
    Elapsed=time.time()-TimeIn
214
    print("Allocation on Device for results : %.3f" % Elapsed)
215

  
216
    TimeIn=time.time()
217
    # Synthesis of function "sillysum" inside Kernel Sources
218
    knl = prg.sillysum  # Use this Kernel object for repeated calls
219
    Elapsed=time.time()-TimeIn
220
    print("Synthesis of kernel : %.3f" % Elapsed)
221

  
222
    TimeIn=time.time()
223
    # Call of kernel previously defined 
224
    CallCL=knl(queue, a_np.shape, None, a_g, b_g, res_g)
225
    # 
226
    CallCL.wait()
227
    Elapsed=time.time()-TimeIn
228
    print("Execution of kernel : %.3f" % Elapsed)
229

  
230
    TimeIn=time.time()
231
    # Creation of vector for result with same size as input vectors
232
    res_np = np.empty_like(a_np)
233
    Elapsed=time.time()-TimeIn
234
    print("Allocation on Host for results: %.3f" % Elapsed)
235

  
236
    TimeIn=time.time()
237
    # Copy from Device to Host
238
    cl.enqueue_copy(queue, res_np, res_g)
239
    Elapsed=time.time()-TimeIn
240
    print("Copy from Device 2 Host : %.3f" % Elapsed)
241

  
242
    return(res_np)
243

  
244
import sys
245
import time
246

  
247
if __name__=='__main__':
248

  
249
    # Size of input vectors definition based on stdin
250
    import sys
251
    try:
252
        SIZE=int(sys.argv[1])
253
        print("Size of vectors set to %i" % SIZE)
254
    except: 
255
        SIZE=50000
256
        print("Size of vectors set to default size %i" % SIZE)
257
        
258
    # a_np = np.random.rand(SIZE).astype(np.float32)
259
    # b_np = np.random.rand(SIZE).astype(np.float32)
260
    
261
    a_np = np.ones(SIZE).astype(np.float32)
262
    b_np = np.ones(SIZE).astype(np.float32)
263

  
264
    # Native & Naive Implementation
265
    print("Performing naive implementation")
266
    TimeIn=time.time()
267
    c_np,d_np=MyDFT(a_np,b_np)
268
    NativeElapsed=time.time()-TimeIn
269
    NativeRate=int(SIZE/NativeElapsed)
270
    print("NativeRate: %i" % NativeRate)
271
    print(c_np,d_np)
272
    
273
    # Native & Numpy Implementation
274
    print("Performing Numpy implementation")
275
    TimeIn=time.time()
276
    e_np,f_np=NumpyDFT(a_np,b_np)
277
    NumpyElapsed=time.time()-TimeIn
278
    NumpyRate=int(SIZE/NumpyElapsed)
279
    print("NumpyRate: %i" % NumpyRate)
280
    print(e_np,f_np) 
281

  
282
    print(np.linalg.norm(c_np-e_np))
283
    print(np.linalg.norm(d_np-f_np))
284
    
285
   #  # OpenCL Implementation
286
   #  TimeIn=time.time()
287
   #  # res_cl=OpenCLAddition(a_np,b_np)
288
   #  res_cl=OpenCLSillyAddition(a_np,b_np)
289
   #  OpenCLElapsed=time.time()-TimeIn
290
   #  OpenCLRate=int(SIZE/OpenCLElapsed)
291
   #  print("OpenCLRate: %i" % OpenCLRate)
292

  
293
   #  # CUDA Implementation
294
   #  TimeIn=time.time()
295
   #  # res_cuda=CUDAAddition(a_np,b_np)
296
   #  res_cuda=CUDASillyAddition(a_np,b_np)
297
   #  CUDAElapsed=time.time()-TimeIn
298
   #  CUDARate=int(SIZE/CUDAElapsed)
299
   #  print("CUDARate: %i" % CUDARate)
300
    
301
   #  print("OpenCLvsNative ratio: %f" % (OpenCLRate/NativeRate))
302
   #  print("CUDAvsNative ratio: %f" % (CUDARate/NativeRate))
303
    
304
   # # Check on OpenCL with Numpy:
305
   #  print(res_cl - res_np)
306
   #  print(np.linalg.norm(res_cl - res_np))
307
   #  try:
308
   #      assert np.allclose(res_np, res_cl)
309
   #  except:
310
   #      print("Results between Native & OpenCL seem to be too different!")
311
        
312
   #  # Check on CUDA with Numpy:
313
   #  print(res_cuda - res_np)
314
   #  print(np.linalg.norm(res_cuda - res_np))
315
   #  try:
316
   #      assert np.allclose(res_np, res_cuda)
317
   #  except:
318
   #      print("Results between Native & CUDA seem to be too different!")
319

  
320

  
0 321

  
ETSN/MyDFT_3.py (revision 270)
1
#!/usr/bin/env python3
2

  
3
import numpy as np
4
import pyopencl as cl
5
from numpy import pi,cos,sin
6

  
7
# piling 16 arithmetical functions
8
def MySillyFunction(x):
9
    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))
10

  
11
# Native Operation under Numpy (for prototyping & tests
12
def NativeAddition(a_np,b_np):
13
    return(a_np+b_np)
14

  
15
# Native Operation with MySillyFunction under Numpy (for prototyping & tests
16
def NativeSillyAddition(a_np,b_np):
17
    return(MySillyFunction(a_np)+MySillyFunction(b_np))
18

  
19
# Naive Discrete Fourier Transform
20
def MyDFT(x,y):
21
    from numpy import pi,cos,sin
22
    size=x.shape[0]
23
    X=np.zeros(size).astype(np.float32)
24
    Y=np.zeros(size).astype(np.float32)
25
    for i in range(size):
26
        for j in range(size):
27
            X[i]=X[i]+x[j]*cos(2.*pi*i*j/size)-y[j]*sin(2.*pi*i*j/size)
28
            Y[i]=Y[i]+x[j]*sin(2.*pi*i*j/size)+y[j]*cos(2.*pi*i*j/size)
29
    return(X,Y)
30

  
31
# Numpy Discrete Fourier Transform
32
def NumpyDFT(x,y):
33
    size=x.shape[0]
34
    X=np.zeros(size).astype(np.float32)
35
    Y=np.zeros(size).astype(np.float32)
36
    nj=np.multiply(2.0*np.pi/size,np.arange(size)).astype(np.float32)
37
    for i in range(size):
38
        X[i]=np.sum(np.subtract(np.multiply(np.cos(i*nj),x),np.multiply(np.sin(i*nj),y)))
39
        Y[i]=np.sum(np.add(np.multiply(np.sin(i*nj),x),np.multiply(np.cos(i*nj),y)))
40
    return(X,Y)
41

  
42
# Numba Discrete Fourier Transform
43
import numba
44
@numba.njit(parallel=True)
45
def NumbaDFT(x,y):
46
    size=x.shape[0]
47
    X=np.zeros(size)
48
    Y=np.zeros(size)
49
    nj=np.multiply(2.0*np.pi/size,np.arange(size)).astype(np.float32)
50
    for i in numba.prange(size):
51
        X[i]=np.sum(np.subtract(np.multiply(np.cos(i*nj),x),np.multiply(np.sin(i*nj),y)))
52
        Y[i]=np.sum(np.add(np.multiply(np.sin(i*nj),x),np.multiply(np.cos(i*nj),y)))
53
    return(X,Y)
54

  
55
# CUDA complete operation
56
def CUDAAddition(a_np,b_np):
57
    import pycuda.autoinit
58
    import pycuda.driver as drv
59
    import numpy
60

  
61
    from pycuda.compiler import SourceModule
62
    mod = SourceModule("""
63
    __global__ void sum(float *dest, float *a, float *b)
64
{
65
  // const int i = threadIdx.x;
66
  const int i = blockIdx.x;
67
  dest[i] = a[i] + b[i];
68
}
69
""")
70

  
71
    # sum = mod.get_function("sum")
72
    sum = mod.get_function("sum")
73

  
74
    res_np = numpy.zeros_like(a_np)
75
    sum(drv.Out(res_np), drv.In(a_np), drv.In(b_np),
76
        block=(1,1,1), grid=(a_np.size,1))
77
    return(res_np)
78

  
79
# CUDA Silly complete operation
80
def CUDASillyAddition(a_np,b_np):
81
    import pycuda.autoinit
82
    import pycuda.driver as drv
83
    import numpy
84

  
85
    from pycuda.compiler import SourceModule
86
    TimeIn=time.time()
87
    mod = SourceModule("""
88
__device__ float MySillyFunction(float x)
89
{
90
    return(pow(sqrt(log(exp(atanh(tanh(asinh(sinh(acosh(cosh(atan(tan(asin(sin(acos(cos(x))))))))))))))),2)); 
91
}
92

  
93
__global__ void sillysum(float *dest, float *a, float *b)
94
{
95
  const int i = blockIdx.x;
96
  dest[i] = MySillyFunction(a[i]) + MySillyFunction(b[i]);
97
}
98
""")
99
    Elapsed=time.time()-TimeIn
100
    print("Definition of kernel : %.3f" % Elapsed)
101

  
102
    TimeIn=time.time()
103
    # sum = mod.get_function("sum")
104
    sillysum = mod.get_function("sillysum")
105
    Elapsed=time.time()-TimeIn
106
    print("Synthesis of kernel : %.3f" % Elapsed)
107

  
108
    TimeIn=time.time()
109
    res_np = numpy.zeros_like(a_np)
110
    Elapsed=time.time()-TimeIn
111
    print("Allocation on Host for results : %.3f" % Elapsed)
112

  
113
    TimeIn=time.time()
114
    sillysum(drv.Out(res_np), drv.In(a_np), drv.In(b_np),
115
             block=(1,1,1), grid=(a_np.size,1))
116
    Elapsed=time.time()-TimeIn
117
    print("Execution of kernel : %.3f" % Elapsed)
118
    return(res_np)
119

  
120
# OpenCL complete operation
121
def OpenCLAddition(a_np,b_np):
122

  
123
    # Context creation
124
    ctx = cl.create_some_context()
125
    # Every process is stored in a queue
126
    queue = cl.CommandQueue(ctx)
127

  
128
    TimeIn=time.time()
129
    # Copy from Host to Device using pointers
130
    mf = cl.mem_flags
131
    a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np)
132
    b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np)
133
    Elapsed=time.time()-TimeIn
134
    print("Copy from Host 2 Device : %.3f" % Elapsed)
135

  
136
    TimeIn=time.time()
137
    # Definition of kernel under OpenCL
138
    prg = cl.Program(ctx, """
139
__kernel void sum(
140
    __global const float *a_g, __global const float *b_g, __global float *res_g)
141
{
142
  int gid = get_global_id(0);
143
  res_g[gid] = a_g[gid] + b_g[gid];
144
}
145
""").build()
146
    Elapsed=time.time()-TimeIn
147
    print("Building kernels : %.3f" % Elapsed)
148
    
149
    TimeIn=time.time()
150
    # Memory allocation on Device for result
151
    res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes)
152
    Elapsed=time.time()-TimeIn
153
    print("Allocation on Device for results : %.3f" % Elapsed)
154

  
155
    TimeIn=time.time()
156
    # Synthesis of function "sum" inside Kernel Sources
157
    knl = prg.sum  # Use this Kernel object for repeated calls
158
    Elapsed=time.time()-TimeIn
159
    print("Synthesis of kernel : %.3f" % Elapsed)
160

  
161
    TimeIn=time.time()
162
    # Call of kernel previously defined 
163
    knl(queue, a_np.shape, None, a_g, b_g, res_g)
164
    Elapsed=time.time()-TimeIn
165
    print("Execution of kernel : %.3f" % Elapsed)
166

  
167
    TimeIn=time.time()
168
    # Creation of vector for result with same size as input vectors
169
    res_np = np.empty_like(a_np)
170
    Elapsed=time.time()-TimeIn
171
    print("Allocation on Host for results: %.3f" % Elapsed)
172

  
173
    TimeIn=time.time()
174
    # Copy from Device to Host
175
    cl.enqueue_copy(queue, res_np, res_g)
176
    Elapsed=time.time()-TimeIn
177
    print("Copy from Device 2 Host : %.3f" % Elapsed)
178

  
179
    return(res_np)
180

  
181
# OpenCL complete operation
182
def OpenCLSillyAddition(a_np,b_np):
183

  
184
    # Context creation
185
    ctx = cl.create_some_context()
186
    # Every process is stored in a queue
187
    queue = cl.CommandQueue(ctx)
188

  
189
    TimeIn=time.time()
190
    # Copy from Host to Device using pointers
191
    mf = cl.mem_flags
192
    a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np)
193
    b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np)
194
    Elapsed=time.time()-TimeIn
195
    print("Copy from Host 2 Device : %.3f" % Elapsed)
196

  
197
    TimeIn=time.time()
198
    # Definition of kernel under OpenCL
199
    prg = cl.Program(ctx, """
200

  
201
float MySillyFunction(float x)
202
{
203
    return(pow(sqrt(log(exp(atanh(tanh(asinh(sinh(acosh(cosh(atan(tan(asin(sin(acos(cos(x))))))))))))))),2)); 
204
}
205

  
206
__kernel void sillysum(
207
    __global const float *a_g, __global const float *b_g, __global float *res_g)
208
{
209
  int gid = get_global_id(0);
210
  res_g[gid] = MySillyFunction(a_g[gid]) + MySillyFunction(b_g[gid]);
211
}
212

  
213
__kernel void sum(
214
    __global const float *a_g, __global const float *b_g, __global float *res_g)
215
{
216
  int gid = get_global_id(0);
217
  res_g[gid] = a_g[gid] + b_g[gid];
218
}
219
""").build()
220
    Elapsed=time.time()-TimeIn
221
    print("Building kernels : %.3f" % Elapsed)
222
    
223
    TimeIn=time.time()
224
    # Memory allocation on Device for result
225
    res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes)
226
    Elapsed=time.time()-TimeIn
227
    print("Allocation on Device for results : %.3f" % Elapsed)
228

  
229
    TimeIn=time.time()
230
    # Synthesis of function "sillysum" inside Kernel Sources
231
    knl = prg.sillysum  # Use this Kernel object for repeated calls
232
    Elapsed=time.time()-TimeIn
233
    print("Synthesis of kernel : %.3f" % Elapsed)
234

  
235
    TimeIn=time.time()
236
    # Call of kernel previously defined 
237
    CallCL=knl(queue, a_np.shape, None, a_g, b_g, res_g)
238
    # 
239
    CallCL.wait()
240
    Elapsed=time.time()-TimeIn
241
    print("Execution of kernel : %.3f" % Elapsed)
242

  
243
    TimeIn=time.time()
244
    # Creation of vector for result with same size as input vectors
245
    res_np = np.empty_like(a_np)
246
    Elapsed=time.time()-TimeIn
247
    print("Allocation on Host for results: %.3f" % Elapsed)
248

  
249
    TimeIn=time.time()
250
    # Copy from Device to Host
251
    cl.enqueue_copy(queue, res_np, res_g)
252
    Elapsed=time.time()-TimeIn
253
    print("Copy from Device 2 Host : %.3f" % Elapsed)
254

  
255
    return(res_np)
256

  
257
import sys
258
import time
259

  
260
if __name__=='__main__':
261

  
262
    # Size of input vectors definition based on stdin
263
    import sys
264
    try:
265
        SIZE=int(sys.argv[1])
266
        print("Size of vectors set to %i" % SIZE)
267
    except: 
268
        SIZE=50000
269
        print("Size of vectors set to default size %i" % SIZE)
270
        
271
    # a_np = np.random.rand(SIZE).astype(np.float32)
272
    # b_np = np.random.rand(SIZE).astype(np.float32)
273
    
274
    a_np = np.ones(SIZE).astype(np.float32)
275
    b_np = np.ones(SIZE).astype(np.float32)
276

  
277
    # Native & Naive Implementation
278
    print("Performing naive implementation")
279
    TimeIn=time.time()
280
    c_np,d_np=MyDFT(a_np,b_np)
281
    NativeElapsed=time.time()-TimeIn
282
    NativeRate=int(SIZE/NativeElapsed)
283
    print("NativeRate: %i" % NativeRate)
284
    
285
    # Native & Numpy Implementation
286
    print("Performing Numpy implementation")
287
    TimeIn=time.time()
288
    e_np,f_np=NumpyDFT(a_np,b_np)
289
    NumpyElapsed=time.time()-TimeIn
290
    NumpyRate=int(SIZE/NumpyElapsed)
291
    print("NumpyRate: %i" % NumpyRate)
292

  
293
    print(np.linalg.norm(c_np-e_np))
294
    print(np.linalg.norm(d_np-f_np))
295
    
296
    # Native & Numpy Implementation
297
    print("Performing Numba implementation")
298
    TimeIn=time.time()
299
    g_np,h_np=NumbaDFT(a_np,b_np)
300
    NumpyElapsed=time.time()-TimeIn
301
    NumpyRate=int(SIZE/NumpyElapsed)
302
    print("NumpyRate: %i" % NumpyRate)
303

  
304
    print(np.linalg.norm(c_np-g_np))
305
    print(np.linalg.norm(d_np-h_np))
306
    
307
   #  # OpenCL Implementation
308
   #  TimeIn=time.time()
309
   #  # res_cl=OpenCLAddition(a_np,b_np)
310
   #  res_cl=OpenCLSillyAddition(a_np,b_np)
311
   #  OpenCLElapsed=time.time()-TimeIn
312
   #  OpenCLRate=int(SIZE/OpenCLElapsed)
313
   #  print("OpenCLRate: %i" % OpenCLRate)
314

  
315
   #  # CUDA Implementation
316
   #  TimeIn=time.time()
317
   #  # res_cuda=CUDAAddition(a_np,b_np)
318
   #  res_cuda=CUDASillyAddition(a_np,b_np)
319
   #  CUDAElapsed=time.time()-TimeIn
320
   #  CUDARate=int(SIZE/CUDAElapsed)
321
   #  print("CUDARate: %i" % CUDARate)
322
    
323
   #  print("OpenCLvsNative ratio: %f" % (OpenCLRate/NativeRate))
324
   #  print("CUDAvsNative ratio: %f" % (CUDARate/NativeRate))
325
    
326
   # # Check on OpenCL with Numpy:
327
   #  print(res_cl - res_np)
328
   #  print(np.linalg.norm(res_cl - res_np))
329
   #  try:
330
   #      assert np.allclose(res_np, res_cl)
331
   #  except:
332
   #      print("Results between Native & OpenCL seem to be too different!")
333
        
334
   #  # Check on CUDA with Numpy:
335
   #  print(res_cuda - res_np)
336
   #  print(np.linalg.norm(res_cuda - res_np))
337
   #  try:
338
   #      assert np.allclose(res_np, res_cuda)
339
   #  except:
340
   #      print("Results between Native & CUDA seem to be too different!")
341

  
342

  
0 343

  

Formats disponibles : Unified diff