Newer
Older
#!/usr/bin/env python3
#
# Pi-by-MonteCarlo using PyCUDA/PyOpenCL
#
# performs an estimation of Pi using Monte Carlo method
# a large amount of iterations is divided and distributed to compute units
# a lot of options are provided to perform scalabilty tests
#
# use -h for complete set of options
#
# CC BY-NC-SA 2011 : Emmanuel QUEMENER <emmanuel.quemener@gmail.com>
# Cecill v2 : Emmanuel QUEMENER <emmanuel.quemener@gmail.com>
#
# Thanks to Andreas Klockner for PyCUDA:
# http://mathema.tician.de/software/pycuda
# Thanks to Andreas Klockner for PyOpenCL:
# http://mathema.tician.de/software/pyopencl
# 2013-01-01 : problems with launch timeout
# http://stackoverflow.com/questions/497685/how-do-you-get-around-the-maximum-cuda-run-time
# Option "Interactive" "0" in /etc/X11/xorg.conf
# Common tools
import numpy
import sys
import getopt
import time
import itertools
from socket import gethostname
class PenStacle:
"""Pentacle of Statistics from data"""
Avg = 0
Med = 0
Std = 0
Min = 0
Max = 0
def __init__(self, Data):
self.Avg = numpy.average(Data)
self.Med = numpy.median(Data)
self.Std = numpy.std(Data)
self.Max = numpy.max(Data)
self.Min = numpy.min(Data)
print("%s %s %s %s %s" % (self.Avg, self.Med, self.Std, self.Min, self.Max))
DeviceStyle = ""
DeviceId = 0
AvgD = 0
MedD = 0
StdD = 0
MinD = 0
MaxD = 0
AvgR = 0
MedR = 0
StdR = 0
MinR = 0
MaxR = 0
def __init__(self, DeviceStyle, DeviceId, Iterations):
self.DeviceStyle = DeviceStyle
self.DeviceId = DeviceId
def Metrology(self, Data):
Duration = PenStacle(Data)
Rate = PenStacle(Iterations / Data)
print("Duration %s" % Duration)
print("Rate %s" % Rate)
def DictionariesAPI():
Marsaglia = {"CONG": 0, "SHR3": 1, "MWC": 2, "KISS": 3}
Computing = {"INT32": 0, "INT64": 1, "FP32": 2, "FP64": 3}
Test = {True: 1, False: 0}
return (Marsaglia, Computing, Test)
# find prime factors of a number
# Get for WWW :
# http://pythonism.wordpress.com/2008/05/17/looking-at-factorisation-in-python/
def PrimeFactors(x):
factorlist = numpy.array([]).astype("uint32")
loop = 2
while loop <= x:
if x % loop == 0:
x /= loop
factorlist = numpy.append(factorlist, [loop])
# Try to find the best thread number in Hybrid approach (Blocks&Threads)
# output is thread number
def BestThreadsNumber(jobs):
factors = PrimeFactors(jobs)
matrix = numpy.append([factors], [factors[::-1]], axis=0)
threads = 1
threads = threads * factor
if threads * threads > jobs or threads > 512:
# Predicted Amdahl Law (Reduced with s=1-p)
# Predicted Amdahl Law
def Amdahl(N, T1, s, p):
def Mylq(N, T1, s, c, p):
return T1 * (s + p / N) + c * N
def Mylq2(N, T1, s, c1, c2, p):
return T1 * (s + p / N) + c1 * N + c2 * N * N
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
#define TCONG 0
#define TSHR3 1
#define TMWC 2
#define TKISS 3
#define TINT32 0
#define TINT64 1
#define TFP32 2
#define TFP64 3
#define IFTHEN 1
// Marsaglia RNG very simple implementation
#define znew ((z=36969*(z&65535)+(z>>16))<<16)
#define wnew ((w=18000*(w&65535)+(w>>16))&65535)
#define MWC (znew+wnew)
#define SHR3 (jsr=(jsr=(jsr=jsr^(jsr<<17))^(jsr>>13))^(jsr<<5))
#define CONG (jcong=69069*jcong+1234567)
#define KISS ((MWC^CONG)+SHR3)
#define MWCfp MWC * 2.328306435454494e-10f
#define KISSfp KISS * 2.328306435454494e-10f
#define SHR3fp SHR3 * 2.328306435454494e-10f
#define CONGfp CONG * 2.328306435454494e-10f
__device__ ulong MainLoop(ulong iterations,uint seed_w,uint seed_z,size_t work)
{
#if TRNG == TCONG
uint jcong=seed_z+work;
#elif TRNG == TSHR3
uint jsr=seed_w+work;
#elif TRNG == TMWC
uint z=seed_z+work;
uint w=seed_w+work;
#elif TRNG == TKISS
uint jcong=seed_z+work;
uint jsr=seed_w+work;
uint z=seed_z-work;
uint w=seed_w-work;
#endif
ulong total=0;
for (ulong i=0;i<iterations;i++) {
#if TYPE == TINT32
#define THEONE 1073741824
#if TRNG == TCONG
uint x=CONG>>17 ;
uint y=CONG>>17 ;
#elif TRNG == TSHR3
uint x=SHR3>>17 ;
uint y=SHR3>>17 ;
#elif TRNG == TMWC
uint x=MWC>>17 ;
uint y=MWC>>17 ;
#elif TRNG == TKISS
uint x=KISS>>17 ;
uint y=KISS>>17 ;
#endif
#elif TYPE == TINT64
#define THEONE 4611686018427387904
#if TRNG == TCONG
ulong x=(ulong)(CONG>>1) ;
ulong y=(ulong)(CONG>>1) ;
#elif TRNG == TSHR3
ulong x=(ulong)(SHR3>>1) ;
ulong y=(ulong)(SHR3>>1) ;
#elif TRNG == TMWC
ulong x=(ulong)(MWC>>1) ;
ulong y=(ulong)(MWC>>1) ;
#elif TRNG == TKISS
ulong x=(ulong)(KISS>>1) ;
ulong y=(ulong)(KISS>>1) ;
#endif
#elif TYPE == TFP32
#define THEONE 1.0f
#if TRNG == TCONG
float x=CONGfp ;
float y=CONGfp ;
#elif TRNG == TSHR3
float x=SHR3fp ;
float y=SHR3fp ;
#elif TRNG == TMWC
float x=MWCfp ;
float y=MWCfp ;
#elif TRNG == TKISS
float x=KISSfp ;
float y=KISSfp ;
#endif
#elif TYPE == TFP64
#define THEONE 1.0f
#if TRNG == TCONG
double x=(double)CONGfp ;
double y=(double)CONGfp ;
#elif TRNG == TSHR3
double x=(double)SHR3fp ;
double y=(double)SHR3fp ;
#elif TRNG == TMWC
double x=(double)MWCfp ;
double y=(double)MWCfp ;
#elif TRNG == TKISS
double x=(double)KISSfp ;
double y=(double)KISSfp ;
#endif
#endif
#if TEST == IFTHEN
if ((x*x+y*y) <=THEONE) {
total+=1;
}
#else
ulong inside=((x*x+y*y) <= THEONE) ? 1:0;
total+=inside;
#endif
}
return(total);
}
__global__ void MainLoopBlocks(ulong *s,ulong iterations,uint seed_w,uint seed_z)
{
ulong total=MainLoop(iterations,seed_z,seed_w,blockIdx.x);
s[blockIdx.x]=total;
__syncthreads();
}
__global__ void MainLoopThreads(ulong *s,ulong iterations,uint seed_w,uint seed_z)
{
ulong total=MainLoop(iterations,seed_z,seed_w,threadIdx.x);
s[threadIdx.x]=total;
__syncthreads();
}
__global__ void MainLoopHybrid(ulong *s,ulong iterations,uint seed_w,uint seed_z)
{
ulong total=MainLoop(iterations,seed_z,seed_w,blockDim.x*blockIdx.x+threadIdx.x);
s[blockDim.x*blockIdx.x+threadIdx.x]=total;
__syncthreads();
}
"""
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
343
344
345
346
347
348
349
350
351
352
353
354
355
356
357
358
359
360
361
362
363
364
365
366
367
368
369
370
371
372
373
374
375
376
377
378
379
380
381
382
383
384
385
386
387
388
389
390
391
392
393
394
395
396
397
398
399
400
401
402
403
404
405
406
407
408
409
410
#define TCONG 0
#define TSHR3 1
#define TMWC 2
#define TKISS 3
#define TINT32 0
#define TINT64 1
#define TFP32 2
#define TFP64 3
#define IFTHEN 1
// Marsaglia RNG very simple implementation
#define znew ((z=36969*(z&65535)+(z>>16))<<16)
#define wnew ((w=18000*(w&65535)+(w>>16))&65535)
#define MWC (znew+wnew)
#define SHR3 (jsr=(jsr=(jsr=jsr^(jsr<<17))^(jsr>>13))^(jsr<<5))
#define CONG (jcong=69069*jcong+1234567)
#define KISS ((MWC^CONG)+SHR3)
#define MWCfp MWC * 2.328306435454494e-10f
#define KISSfp KISS * 2.328306435454494e-10f
#define CONGfp CONG * 2.328306435454494e-10f
#define SHR3fp SHR3 * 2.328306435454494e-10f
ulong MainLoop(ulong iterations,uint seed_z,uint seed_w,size_t work)
{
#if TRNG == TCONG
uint jcong=seed_z+work;
#elif TRNG == TSHR3
uint jsr=seed_w+work;
#elif TRNG == TMWC
uint z=seed_z+work;
uint w=seed_w+work;
#elif TRNG == TKISS
uint jcong=seed_z+work;
uint jsr=seed_w+work;
uint z=seed_z-work;
uint w=seed_w-work;
#endif
ulong total=0;
for (ulong i=0;i<iterations;i++) {
#if TYPE == TINT32
#define THEONE 1073741824
#if TRNG == TCONG
uint x=CONG>>17 ;
uint y=CONG>>17 ;
#elif TRNG == TSHR3
uint x=SHR3>>17 ;
uint y=SHR3>>17 ;
#elif TRNG == TMWC
uint x=MWC>>17 ;
uint y=MWC>>17 ;
#elif TRNG == TKISS
uint x=KISS>>17 ;
uint y=KISS>>17 ;
#endif
#elif TYPE == TINT64
#define THEONE 4611686018427387904
#if TRNG == TCONG
ulong x=(ulong)(CONG>>1) ;
ulong y=(ulong)(CONG>>1) ;
#elif TRNG == TSHR3
ulong x=(ulong)(SHR3>>1) ;
ulong y=(ulong)(SHR3>>1) ;
#elif TRNG == TMWC
ulong x=(ulong)(MWC>>1) ;
ulong y=(ulong)(MWC>>1) ;
#elif TRNG == TKISS
ulong x=(ulong)(KISS>>1) ;
ulong y=(ulong)(KISS>>1) ;
#endif
#elif TYPE == TFP32
#define THEONE 1.0f
#if TRNG == TCONG
float x=CONGfp ;
float y=CONGfp ;
#elif TRNG == TSHR3
float x=SHR3fp ;
float y=SHR3fp ;
#elif TRNG == TMWC
float x=MWCfp ;
float y=MWCfp ;
#elif TRNG == TKISS
float x=KISSfp ;
float y=KISSfp ;
#endif
#elif TYPE == TFP64
#pragma OPENCL EXTENSION cl_khr_fp64: enable
#define THEONE 1.0f
#if TRNG == TCONG
double x=(double)CONGfp ;
double y=(double)CONGfp ;
#elif TRNG == TSHR3
double x=(double)SHR3fp ;
double y=(double)SHR3fp ;
#elif TRNG == TMWC
double x=(double)MWCfp ;
double y=(double)MWCfp ;
#elif TRNG == TKISS
double x=(double)KISSfp ;
double y=(double)KISSfp ;
#endif
#endif
#if TEST == IFTHEN
if ((x*x+y*y) <= THEONE) {
total+=1;
}
#else
ulong inside=((x*x+y*y) <= THEONE) ? 1:0;
total+=inside;
#endif
}
__kernel void MainLoopGlobal(
__global ulong *s,ulong iterations,uint seed_w,uint seed_z)
{
ulong total=MainLoop(iterations,seed_z,seed_w,get_global_id(0));
barrier(CLK_GLOBAL_MEM_FENCE);
__kernel void MainLoopLocal(
__global ulong *s,ulong iterations,uint seed_w,uint seed_z)
{
ulong total=MainLoop(iterations,seed_z,seed_w,get_local_id(0));
barrier(CLK_LOCAL_MEM_FENCE);
s[get_local_id(0)]=total;
}
__kernel void MainLoopHybrid(
__global ulong *s,ulong iterations,uint seed_w,uint seed_z)
{
ulong total=MainLoop(iterations,seed_z,seed_w,get_global_id(0));
barrier(CLK_GLOBAL_MEM_FENCE || CLK_LOCAL_MEM_FENCE);
s[get_global_id(0)]=total;
}
"""
print("Inside ", InputCU)
iterations = InputCU["Iterations"]
steps = InputCU["Steps"]
blocks = InputCU["Blocks"]
threads = InputCU["Threads"]
Device = InputCU["Device"]
RNG = InputCU["RNG"]
ValueType = InputCU["ValueType"]
Seeds = InputCU["Seeds"]
Marsaglia, Computing, Test = DictionariesAPI()
try:
# For PyCUDA import
import pycuda.driver as cuda
from pycuda.compiler import SourceModule
print("GPU selected %s" % XPU.name())
print
except ImportError:
print("Platform does not seem to support CUDA")
circle = numpy.zeros(blocks * threads).astype(numpy.uint64)
# circleCU = cuda.mem_alloc(circle.size*circle.dtype.itemize)
# cuda.memcpy_htod(circleCU, circle)
mod = SourceModule(
KernelCodeCuda(),
options=[
"--compiler-options",
"-DTRNG=%i -DTYPE=%s" % (Marsaglia[RNG], Computing[ValueType]),
],
)
# mod = SourceModule(KernelCodeCuda(),nvcc='nvcc',keep=True)
# Needed to set the compiler via ccbin for CUDA9 implementation
# mod = SourceModule(KernelCodeCuda(),options=['-ccbin','clang-3.9','--compiler-options','-DTRNG=%i' % Marsaglia[RNG],'-DTYPE=%s' % Computing[ValueType],'-DTEST=%s' % Test[TestType]],keep=True) # noqa: E501
except Exception:
MetropolisBlocksCU = mod.get_function("MainLoopBlocks") # noqa: F841
MetropolisThreadsCU = mod.get_function("MainLoopThreads") # noqa: F841
MetropolisHybridCU = mod.get_function("MainLoopHybrid")
MyDuration = numpy.zeros(steps)
iterationsCU = numpy.uint64(iterations / jobs)
if iterations % jobs != 0:
iterationsCU += numpy.uint64(1)
MetropolisHybridCU(
circleCU,
numpy.uint64(iterationsCU),
numpy.uint32(Seeds[0]),
numpy.uint32(Seeds[1]),
grid=(blocks, 1),
block=(threads, 1, 1),
)
except Exception:
elapsed = time.time() - start_time
print(
"(Blocks/Threads)=(%i,%i) method done in %.2f s..."
% (blocks, threads, elapsed)
)
OutputCU = {
"Inside": sum(circle),
"NewIterations": numpy.uint64(iterationsCU * jobs),
"Duration": MyDuration,
}
def MetropolisOpenCL(InputCL):
import pyopencl as cl
iterations = InputCL["Iterations"]
steps = InputCL["Steps"]
blocks = InputCL["Blocks"]
threads = InputCL["Threads"]
Device = InputCL["Device"]
RNG = InputCL["RNG"]
ValueType = InputCL["ValueType"]
TestType = InputCL["IfThen"]
Seeds = InputCL["Seeds"]
Marsaglia, Computing, Test = DictionariesAPI()
# Initialisation des variables en les CASTant correctement
for platform in cl.get_platforms():
for device in platform.get_devices():
if Id == Device:
XPU = device
print("CPU/GPU selected: ", device.name.lstrip())
HasXPU = True
Id += 1
if not HasXPU:
print("No XPU #%i found in all of %i devices, sorry..." % (Device, Id - 1))
sys.exit()
# Je cree le contexte et la queue pour son execution
try:
ctx = cl.Context(devices=[XPU])
queue = cl.CommandQueue(
ctx, properties=cl.command_queue_properties.PROFILING_ENABLE
)
except Exception:
print("Crash during context creation")
# Je recupere les flag possibles pour les buffers
mf = cl.mem_flags
circle = numpy.zeros(blocks * threads).astype(numpy.uint64)
circleCL = cl.Buffer(ctx, mf.WRITE_ONLY | mf.COPY_HOST_PTR, hostbuf=circle)
MetropolisCL = cl.Program(ctx, KernelCodeOpenCL()).build(
options="-cl-mad-enable -cl-fast-relaxed-math -DTRNG=%i -DTYPE=%s -DTEST=%s"
% (Marsaglia[RNG], Computing[ValueType], Test[TestType])
)
MyDuration = numpy.zeros(steps)
jobs = blocks * threads
iterationsCL = numpy.uint64(iterations / jobs)
if iterations % jobs != 0:
iterationsCL += 1
CLLaunch = MetropolisCL.MainLoopGlobal(
queue,
(blocks,),
None,
circleCL,
numpy.uint64(iterationsCL),
numpy.uint32(Seeds[0]),
numpy.uint32(Seeds[1]),
)
CLLaunch = MetropolisCL.MainLoopHybrid(
queue,
(jobs,),
(threads,),
circleCL,
numpy.uint64(iterationsCL),
numpy.uint32(Seeds[0]),
numpy.uint32(Seeds[1]),
)
CLLaunch.wait()
cl.enqueue_copy(queue, circle, circleCL).wait()
elapsed = time.time() - start_time
print(
"(Blocks/Threads)=(%i,%i) method done in %.2f s..."
% (blocks, threads, elapsed)
)
# Elapsed method based on CLLaunch doesn't work for Beignet OpenCL
# elapsed = 1e-9*(CLLaunch.profile.end - CLLaunch.profile.start)
# print circle,numpy.mean(circle),numpy.median(circle),numpy.std(circle)
# AllPi=4./numpy.float32(iterationsCL)*circle.astype(numpy.float32)
# MyPi[i]=numpy.median(AllPi)
# print MyPi[i],numpy.std(AllPi),MyDuration[i]
circleCL.release()
OutputCL = {
"Inside": sum(circle),
"NewIterations": numpy.uint64(iterationsCL * jobs),
"Duration": MyDuration,
}
from scipy.optimize import curve_fit
import matplotlib.pyplot as plt
try:
coeffs_Amdahl, matcov_Amdahl = curve_fit(Amdahl, N, D)
D_Amdahl = Amdahl(N, coeffs_Amdahl[0], coeffs_Amdahl[1], coeffs_Amdahl[2])
coeffs_Amdahl[1] = coeffs_Amdahl[1] * coeffs_Amdahl[0] / D[0]
coeffs_Amdahl[2] = coeffs_Amdahl[2] * coeffs_Amdahl[0] / D[0]
coeffs_Amdahl[0] = D[0]
print(
"Amdahl Normalized: T=%.2f(%.6f+%.6f/N)"
% (coeffs_Amdahl[0], coeffs_Amdahl[1], coeffs_Amdahl[2])
)
except Exception:
print("Impossible to fit for Amdahl law : only %i elements" % len(D))
try:
coeffs_AmdahlR, matcov_AmdahlR = curve_fit(AmdahlR, N, D)
# D_AmdahlR = AmdahlR(N, coeffs_AmdahlR[0], coeffs_AmdahlR[1])
coeffs_AmdahlR[1] = coeffs_AmdahlR[1] * coeffs_AmdahlR[0] / D[0]
coeffs_AmdahlR[0] = D[0]
print(
"Amdahl Reduced Normalized: T=%.2f(%.6f+%.6f/N)"
% (coeffs_AmdahlR[0], 1 - coeffs_AmdahlR[1], coeffs_AmdahlR[1])
)
print("Impossible to fit for Reduced Amdahl law : only %i elements" % len(D))
try:
coeffs_Mylq, matcov_Mylq = curve_fit(Mylq, N, D)
coeffs_Mylq[1] = coeffs_Mylq[1] * coeffs_Mylq[0] / D[0]
coeffs_Mylq[3] = coeffs_Mylq[3] * coeffs_Mylq[0] / D[0]
coeffs_Mylq[0] = D[0]
print(
"Mylq Normalized : T=%.2f(%.6f+%.6f/N)+%.6f*N"
% (coeffs_Mylq[0], coeffs_Mylq[1], coeffs_Mylq[3], coeffs_Mylq[2])
)
D_Mylq = Mylq(N, coeffs_Mylq[0], coeffs_Mylq[1], coeffs_Mylq[2],
coeffs_Mylq[3])
except Exception:
print("Impossible to fit for Mylq law : only %i elements" % len(D))
try:
coeffs_Mylq2, matcov_Mylq2 = curve_fit(Mylq2, N, D)
coeffs_Mylq2[1] = coeffs_Mylq2[1] * coeffs_Mylq2[0] / D[0]
# coeffs_Mylq2[2]=coeffs_Mylq2[2]*coeffs_Mylq2[0]/D[0]
# coeffs_Mylq2[3]=coeffs_Mylq2[3]*coeffs_Mylq2[0]/D[0]
coeffs_Mylq2[4] = coeffs_Mylq2[4] * coeffs_Mylq2[0] / D[0]
coeffs_Mylq2[0] = D[0]
print(
"Mylq 2nd order Normalized: T=%.2f(%.6f+%.6f/N)+%.6f*N+%.6f*N^2"
% (
coeffs_Mylq2[0],
coeffs_Mylq2[1],
coeffs_Mylq2[4],
coeffs_Mylq2[2],
coeffs_Mylq2[3],
)
)
except Exception:
print("Impossible to fit for 2nd order Mylq law : only %i elements" % len(D))
if Curves:
plt.xlabel("Number of Threads/work Items")
plt.ylabel("Total Elapsed Time")
(pAmdahl,) = plt.plot(N, D_Amdahl, label="Loi de Amdahl")
(pMylq,) = plt.plot(N, D_Mylq, label="Loi de Mylq")
except Exception:
print("Fit curves seem not to be available")
plt.legend()
plt.show()
# GPU style can be Cuda (Nvidia implementation) or OpenCL
# Redo is the times to redo the test to improve metrology
# OutMetrology is method for duration estimation : False is GPU inside
OutMetrology = False
Metrology = "InMetro"
Seeds = 110271, 101008
HowToUse = "%s -o (Out of Core Metrology) -c (Print Curves) -k (Case On IfThen) -d <DeviceId> -g <CUDA/OpenCL> -i <Iterations> -b <BlocksBegin> -e <BlocksEnd> -s <BlocksStep> -f <ThreadsFirst> -l <ThreadsLast> -t <ThreadssTep> -r <RedoToImproveStats> -m <SHR3/CONG/MWC/KISS> -v <INT32/INT64/FP32/FP64>" # noqa: E501
opts, args = getopt.getopt(
sys.argv[1:],
"hockg:i:b:e:s:f:l:t:r:d:m:v:",
[
"gpustyle=",
"iterations=",
"blocksBegin=",
"blocksEnd=",
"blocksStep=",
"threadsFirst=",
"threadsLast=",
"threadssTep=",
"redo=",
"device=",
"marsaglia=",
"valuetype=",
],
)
except getopt.GetoptError:
print(HowToUse % sys.argv[0])
sys.exit(2)
# List of Devices
print(HowToUse % sys.argv[0])
print("\nInformations about devices detected under OpenCL API:")
# For PyOpenCL import
try:
import pyopencl as cl
for platform in cl.get_platforms():
for device in platform.get_devices():
# deviceType=cl.device_type.to_string(device.type)
deviceType = "xPU"
print(
"Device #%i from %s of type %s : %s"
% (
Id,
platform.vendor.lstrip(),
deviceType,
device.name.lstrip(),
)
)
Id = Id + 1
except Exception:
print("Your platform does not seem to support OpenCL")
print("\nInformations about devices detected under CUDA API:")
# For PyCUDA import
try:
import pycuda.driver as cuda
device = cuda.Device(Id)
print("Device #%i of type GPU : %s" % (Id, device.name()))
elif opt == "-o":
OutMetrology = True
Metrology = "OutMetro"
elif opt == "-c":
Curves = True
elif opt == "-k":
IfThen = True
elif opt in ("-d", "--device"):
Devices.append(int(arg))
elif opt in ("-g", "--gpustyle"):
GpuStyle = arg
elif opt in ("-m", "--marsaglia"):
RNG = arg
elif opt in ("-v", "--valuetype"):
ValueType = arg
elif opt in ("-i", "--iterations"):
Iterations = numpy.uint64(arg)
elif opt in ("-b", "--blocksbegin"):
BlocksBegin = int(arg)
elif opt in ("-e", "--blocksend"):
BlocksEnd = int(arg)
elif opt in ("-s", "--blocksstep"):
BlocksStep = int(arg)
elif opt in ("-f", "--threadsfirst"):
ThreadsBegin = int(arg)
ThreadsEnd = ThreadsBegin
elif opt in ("-l", "--threadslast"):
ThreadsEnd = int(arg)
elif opt in ("-t", "--threadsstep"):
ThreadsStep = int(arg)
elif opt in ("-r", "--redo"):
Redo = int(arg)
# If no device has been specified, take the first one!
print("Devices Identification : %s" % Devices)
print("GpuStyle used : %s" % GpuStyle)
print("Iterations : %s" % Iterations)
print("Number of Blocks on begin : %s" % BlocksBegin)
print("Number of Blocks on end : %s" % BlocksEnd)
print("Step on Blocks : %s" % BlocksStep)
print("Number of Threads on begin : %s" % ThreadsBegin)
print("Number of Threads on end : %s" % ThreadsEnd)
print("Step on Threads : %s" % ThreadsStep)
print("Number of redo : %s" % Redo)
print("Metrology done out of XPU : %r" % OutMetrology)
print("Type of Marsaglia RNG used : %s" % RNG)
print("Type of variable : %s" % ValueType)
try:
# For PyCUDA import
import pycuda.driver as cuda
device = cuda.Device(Id)
print("Device #%i of type GPU : %s" % (Id, device.name()))
except ImportError:
print("Platform does not seem to support CUDA")
try:
# For PyOpenCL import
import pyopencl as cl
for platform in cl.get_platforms():
for device in platform.get_devices():
# deviceType=cl.device_type.to_string(device.type)
deviceType = "xPU"
print(
"Device #%i from %s of type %s : %s"
% (
Id,
platform.vendor.lstrip().rstrip(),
deviceType,
device.name.lstrip().rstrip(),
)
)
# Set the Alu as detected Device Type
Alu[Id] = deviceType
Id = Id + 1
except ImportError:
print("Platform does not seem to support OpenCL")
# print(Devices,Alu)
BlocksList = range(BlocksBegin, BlocksEnd + BlocksStep, BlocksStep)
ThreadsList = range(ThreadsBegin, ThreadsEnd + ThreadsStep, ThreadsStep)
ExploredJobs = numpy.array([]).astype(numpy.uint32)
ExploredBlocks = numpy.array([]).astype(numpy.uint32)
ExploredThreads = numpy.array([]).astype(numpy.uint32)
avgD = numpy.array([]).astype(numpy.float32)
medD = numpy.array([]).astype(numpy.float32)
stdD = numpy.array([]).astype(numpy.float32)
minD = numpy.array([]).astype(numpy.float32)
maxD = numpy.array([]).astype(numpy.float32)
avgR = numpy.array([]).astype(numpy.float32)
medR = numpy.array([]).astype(numpy.float32)
stdR = numpy.array([]).astype(numpy.float32)
minR = numpy.array([]).astype(numpy.float32)
maxR = numpy.array([]).astype(numpy.float32)
for Blocks, Threads in itertools.product(BlocksList, ThreadsList):
circle = numpy.zeros(Blocks * Threads).astype(numpy.uint64)
ExploredJobs = numpy.append(ExploredJobs, Blocks * Threads)
ExploredBlocks = numpy.append(ExploredBlocks, Blocks)
ExploredThreads = numpy.append(ExploredThreads, Threads)
if OutMetrology:
DurationItem = numpy.array([]).astype(numpy.float32)
Duration = numpy.array([]).astype(numpy.float32)
Rate = numpy.array([]).astype(numpy.float32)
start = time.time()
if GpuStyle == "CUDA":
InputCU = {}
InputCU["Iterations"] = Iterations
InputCU["Steps"] = 1
InputCU["Blocks"] = Blocks
InputCU["Threads"] = Threads
InputCU["Device"] = Devices[0]
InputCU["RNG"] = RNG
InputCU["Seeds"] = Seeds
InputCU["ValueType"] = ValueType
InputCU["IfThen"] = IfThen
OutputCU = MetropolisCuda(InputCU)
Inside = OutputCU["Circle"]
NewIterations = OutputCU["NewIterations"]
Duration = OutputCU["Duration"]
except Exception:
print(
"Problem with (%i,%i) // computations on Cuda"
% (Blocks, Threads)
)
elif GpuStyle == "OpenCL":