Centre Blaise Pascal » Bench4GPU

#!/usr/bin/env python

#

# Pi-by-MonteCarlo using PyCUDA/PyOpenCL

#

# CC BY-NC-SA 2011 : <emmanuel.quemener@ens-lyon.fr>

#

# Thanks to Andreas Klockner for PyCUDA:

# http://mathema.tician.de/software/pycuda

#

# 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

from numpy.random import randint as nprnd

import sys

import getopt

import time

import math

from socket import gethostname

# 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])

    else:

      loop+=1

  return factorlist

# 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

  for factor in matrix.transpose().ravel():

    threads=threads*factor

    if threads*threads>jobs:

      break

  return(long(threads))

# Predicted Amdahl Law (Reduced with s=1-p)

def AmdahlR(N, T1, p):

  return (T1*(1-p+p/N))

# Predicted Amdahl Law

def Amdahl(N, T1, s, p):

  return (T1*(s+p/N))

# Predicted Mylq Law with first order

def Mylq(N, T1,s,c,p):

  return (T1*(s+p/N)+c*N)

# Predicted Mylq Law with second order

def Mylq2(N, T1,s,c1,c2,p):

  return (T1*(s+p/N)+c1*N+c2*N*N)

KERNEL_CODE_CUDA="""

// 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

__global__ void MainLoopBlocks(ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(blockIdx.x+1);

   uint w=seed_w/(blockIdx.x+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      float x=MWCfp ;

      float y=MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[blockIdx.x]=total;

   __syncthreads();

}

__global__ void MainLoopThreads(ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(threadIdx.x+1);

   uint w=seed_w/(threadIdx.x+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      float x=MWCfp ;

      float y=MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[threadIdx.x]=total;

   __syncthreads();

}

__global__ void MainLoopHybrid(ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(blockDim.x*blockIdx.x+threadIdx.x+1);

   uint w=seed_w/(blockDim.x*blockIdx.x+threadIdx.x+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      float x=MWCfp ;

      float y=MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[blockDim.x*blockIdx.x+threadIdx.x]=total;

   __syncthreads();

}

__global__ void MainLoopBlocks64(ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(blockIdx.x+1);

   uint w=seed_w/(blockIdx.x+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      double x=(double)MWCfp ;

      double y=(double)MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[blockIdx.x]=total;

   __syncthreads();

}

__global__ void MainLoopThreads64(ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(threadIdx.x+1);

   uint w=seed_w/(threadIdx.x+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      double x=(double)MWCfp ;

      double y=(double)MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[threadIdx.x]=total;

   __syncthreads();

}

__global__ void MainLoopHybrid64(ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(blockDim.x*blockIdx.x+threadIdx.x+1);

   uint w=seed_w/(blockDim.x*blockIdx.x+threadIdx.x+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      double x=(double)MWCfp ;

      double y=(double)MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[blockDim.x*blockIdx.x+threadIdx.x]=total;

   __syncthreads();

}

"""

KERNEL_CODE_OPENCL="""

#pragma OPENCL EXTENSION cl_khr_fp64: enable

// 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

__kernel void MainLoopGlobal(__global ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(get_global_id(0)+1);

   uint w=seed_w/(get_global_id(0)+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      float x=MWCfp ;

      float y=MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[get_global_id(0)]=total;

   barrier(CLK_GLOBAL_MEM_FENCE);

}

__kernel void MainLoopLocal(__global ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(get_local_id(0)+1);

   uint w=seed_w/(get_local_id(0)+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      float x=MWCfp ;

      float y=MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[get_local_id(0)]=total;

   barrier(CLK_LOCAL_MEM_FENCE);

}

__kernel void MainLoopHybrid(__global ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(get_group_id(0)*get_num_groups(0)+get_local_id(0)+1);

   uint w=seed_w/(get_group_id(0)*get_num_groups(0)+get_local_id(0)+1);

   ulong total=0;

   for (uint i=0;i<iterations;i++) {

      float x=MWCfp ;

     float y=MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   barrier(CLK_LOCAL_MEM_FENCE);

   s[get_group_id(0)*get_num_groups(0)+get_local_id(0)]=total;

}

__kernel void MainLoopGlobal64(__global ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(get_global_id(0)+1);

   uint w=seed_w/(get_global_id(0)+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      double x=(double)MWCfp ;

      double y=(double)MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[get_global_id(0)]=total;

   barrier(CLK_GLOBAL_MEM_FENCE);

}

__kernel void MainLoopLocal64(__global ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(get_local_id(0)+1);

   uint w=seed_w/(get_local_id(0)+1);

   ulong total=0;

   for (ulong i=0;i<iterations;i++) {

      double x=(double)MWCfp ;

      double y=(double)MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   s[get_local_id(0)]=total;

   barrier(CLK_LOCAL_MEM_FENCE);

}

__kernel void MainLoopHybrid64(__global ulong *s,ulong iterations,uint seed_w,uint seed_z)

{

   uint z=seed_z/(get_group_id(0)*get_num_groups(0)+get_local_id(0)+1);

   uint w=seed_w/(get_group_id(0)*get_num_groups(0)+get_local_id(0)+1);

   ulong total=0;

   for (uint i=0;i<iterations;i++) {

      double x=(double)MWCfp ;

      double y=(double)MWCfp ;

      // Matching test

      ulong inside=((x*x+y*y) < 1.0f) ? 1:0;

      total+=inside;

   }

   barrier(CLK_LOCAL_MEM_FENCE);

   s[get_group_id(0)*get_num_groups(0)+get_local_id(0)]=total;

}

"""

def MetropolisCuda(circle,iterations,steps,jobs,ParaStyle,DoublePrecision):

  # Avec PyCUDA autoinit, rien a faire !

  circleCU = cuda.InOut(circle)

  mod = SourceModule(KERNEL_CODE_CUDA)

  MetropolisBlocksCU=mod.get_function("MainLoopBlocks")

  MetropolisJobsCU=mod.get_function("MainLoopThreads")

  MetropolisHybridCU=mod.get_function("MainLoopHybrid")

  MetropolisBlocks64CU=mod.get_function("MainLoopBlocks64")

  MetropolisJobs64CU=mod.get_function("MainLoopThreads64")

  MetropolisHybrid64CU=mod.get_function("MainLoopHybrid64")

  start = pycuda.driver.Event()

  stop = pycuda.driver.Event()

  MyPi=numpy.zeros(steps)

  MyDuration=numpy.zeros(steps)

  if iterations%jobs==0:

    iterationsCL=numpy.uint64(iterations/jobs)

    iterationsNew=iterationsCL*jobs

  else:

    iterationsCL=numpy.uint64(iterations/jobs+1)

    iterationsNew=iterations

  for i in range(steps):

    start.record()

    start.synchronize()

    if ParaStyle=='Blocks':

      if DoublePrecision:

        MetropolisBlocksCU(circleCU,

                           numpy.uint64(iterationsCL),

                           numpy.uint32(nprnd(2**30/jobs)),

                           numpy.uint32(nprnd(2**30/jobs)),

                           grid=(jobs,1),

                           block=(1,1,1))

      else:

        MetropolisBlocks64CU(circleCU,

                             numpy.uint64(iterationsCL),

                             numpy.uint32(nprnd(2**30/jobs)),

                             numpy.uint32(nprnd(2**30/jobs)),

                             grid=(jobs,1),

                             block=(1,1,1))

      print "%s with (WorkItems/Threads)=(%i,%i) %s method done" % \

            (Alu,jobs,1,ParaStyle)

    elif ParaStyle=='Hybrid':

      threads=BestThreadsNumber(jobs)

      if DoublePrecision:

        MetropolisHybrid64CU(circleCU,

                             numpy.uint64(iterationsCL),

                             numpy.uint32(nprnd(2**30/jobs)),

                             numpy.uint32(nprnd(2**30/jobs)),

                             grid=(jobs,1),

                             block=(threads,1,1))

      else:

        MetropolisHybridCU(circleCU,

                           numpy.uint64(iterationsCL),

                           numpy.uint32(nprnd(2**30/jobs)),

                           numpy.uint32(nprnd(2**30/jobs)),

                           grid=(jobs,1),

                           block=(threads,1,1))

      print "%s with (WorkItems/Threads)=(%i,%i) %s method done" % \

            (Alu,jobs/threads,threads,ParaStyle)

    else:

      if DoublePrecision:

        MetropolisJobs64CU(circleCU,

                           numpy.uint64(iterationsCL),

                           numpy.uint32(nprnd(2**30/jobs)),

                           numpy.uint32(nprnd(2**30/jobs)),

                           grid=(1,1),

                           block=(jobs,1,1))

      else:

        MetropolisJobsCU(circleCU,

                         numpy.uint64(iterationsCL),

                         numpy.uint32(nprnd(2**30/jobs)),

                         numpy.uint32(nprnd(2**30/jobs)),

                         grid=(1,1),

                         block=(jobs,1,1))

      print "%s with (WorkItems/Threads)=(%i,%i) %s method done" % \

            (Alu,jobs,1,ParaStyle)

    stop.record()

    stop.synchronize()

    elapsed = start.time_till(stop)*1e-3

    MyDuration[i]=elapsed

    AllPi=4./numpy.float32(iterationsCL)*circle.astype(numpy.float32)

    MyPi[i]=numpy.median(AllPi)

    print MyPi[i],numpy.std(AllPi),MyDuration[i]

  print jobs,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)

  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))

def MetropolisOpenCL(circle,iterations,steps,jobs,ParaStyle,Alu,Device,

                     DoublePrecision):

  # Initialisation des variables en les CASTant correctement

  if Device==0:

    print "Enter XPU selector based on ALU type: first selected"

    HasXPU=False

    # Default Device selection based on ALU Type

    for platform in cl.get_platforms():

      for device in platform.get_devices():

        deviceType=cl.device_type.to_string(device.type)

        if deviceType=="GPU" and Alu=="GPU" and not HasXPU:

          XPU=device

          print "GPU selected: ",device.name

          HasXPU=True

        if deviceType=="CPU" and Alu=="CPU" and not HasXPU:

          XPU=device

          print "CPU selected: ",device.name

          HasXPU=True

  else:

    print "Enter XPU selector based on device number & ALU type"

    Id=1

    HasXPU=False

    # Primary Device selection based on Device Id

    for platform in cl.get_platforms():

      for device in platform.get_devices():

        deviceType=cl.device_type.to_string(device.type)

        if Id==Device and Alu==deviceType and HasXPU==False:

          XPU=device

          print "CPU/GPU selected: ",device.name

          HasXPU=True

        Id=Id+1

    if HasXPU==False:

      print "No XPU #%i of type %s found in all of %i devices, sorry..." % \

          (Device,Alu,Id-1)

      return(0,0,0)

  # Je cree le contexte et la queue pour son execution

  ctx = cl.Context([XPU])

  queue = cl.CommandQueue(ctx,

                          properties=cl.command_queue_properties.PROFILING_ENABLE)

  # Je recupere les flag possibles pour les buffers

  mf = cl.mem_flags

  circleCL = cl.Buffer(ctx, mf.WRITE_ONLY|mf.COPY_HOST_PTR,hostbuf=circle)

  MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build( \

    options = "-cl-mad-enable -cl-fast-relaxed-math")

  i=0

  MyPi=numpy.zeros(steps)

  MyDuration=numpy.zeros(steps)

  if iterations%jobs==0:

    iterationsCL=numpy.uint64(iterations/jobs)

    iterationsNew=numpy.uint64(iterationsCL*jobs)

  else:

    iterationsCL=numpy.uint64(iterations/jobs+1)

    iterationsNew=numpy.uint64(iterations)

  for i in range(steps):

    if ParaStyle=='Blocks':

      # Call OpenCL kernel

      # (1,) is Global work size (only 1 work size)

      # (1,) is local work size

      # circleCL is lattice translated in CL format

      # SeedZCL is lattice translated in CL format

      # SeedWCL is lattice translated in CL format

      # step is number of iterations

      if DoublePrecision:

        CLLaunch=MetropolisCL.MainLoopGlobal64(queue,(jobs,),None,

                                               circleCL,

                                               numpy.uint64(iterationsCL),

                                               numpy.uint32(nprnd(2**30/jobs)),

                                               numpy.uint32(nprnd(2**30/jobs)))

      else:

        CLLaunch=MetropolisCL.MainLoopGlobal(queue,(jobs,),None,

                                             circleCL,

                                             numpy.uint64(iterationsCL),

                                             numpy.uint32(nprnd(2**30/jobs)),

                                             numpy.uint32(nprnd(2**30/jobs)))

      print "%s with (WorkItems/Threads)=(%i,%i) %s method done" % \

            (Alu,jobs,1,ParaStyle)

    elif ParaStyle=='Hybrid':

      threads=BestThreadsNumber(jobs)

      # en OpenCL, necessaire de mettre un Global_id identique au local_id

      if DoublePrecision:

        CLLaunch=MetropolisCL.MainLoopHybrid64(queue,(jobs,),(threads,),

                                               circleCL,

                                               numpy.uint64(iterationsCL),

                                               numpy.uint32(nprnd(2**30/jobs)),

                                               numpy.uint32(nprnd(2**30/jobs)))

      else:

        CLLaunch=MetropolisCL.MainLoopHybrid(queue,(jobs,),(threads,),

                                             circleCL,

                                             numpy.uint64(iterationsCL),

                                             numpy.uint32(nprnd(2**30/jobs)),

                                             numpy.uint32(nprnd(2**30/jobs)))

      print "%s with (WorkItems/Threads)=(%i,%i) %s method done" % \

            (Alu,jobs/threads,threads,ParaStyle)

    else:

      # en OpenCL, necessaire de mettre un Global_id identique au local_id

      if DoublePrecision:

        CLLaunch=MetropolisCL.MainLoopLocal64(queue,(jobs,),(jobs,),

                                              circleCL,

                                              numpy.uint64(iterationsCL),

                                              numpy.uint32(nprnd(2**30/jobs)),

                                              numpy.uint32(nprnd(2**30/jobs)))

      else:

        CLLaunch=MetropolisCL.MainLoopLocal(queue,(jobs,),(jobs,),

                                            circleCL,

                                            numpy.uint64(iterationsCL),

                                            numpy.uint32(nprnd(2**30/jobs)),

                                            numpy.uint32(nprnd(2**30/jobs)))

      print "%s with %i %s done" % (Alu,jobs,ParaStyle)

    CLLaunch.wait()

    cl.enqueue_copy(queue, circle, circleCL).wait()

    elapsed = 1e-9*(CLLaunch.profile.end - CLLaunch.profile.start)

    MyDuration[i]=elapsed

    AllPi=4./numpy.float32(iterationsCL)*circle.astype(numpy.float32)

    MyPi[i]=numpy.median(AllPi)

    print MyPi[i],numpy.std(AllPi),MyDuration[i]

  circleCL.release()

  print jobs,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)

  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))

def FitAndPrint(N,D,Curves):

  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:

    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])

  except:

    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[2]=coeffs_Mylq[2]*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:

    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:

    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")

    Experience,=plt.plot(N,D,'ro')

    try:

      pAmdahl,=plt.plot(N,D_Amdahl,label="Loi de Amdahl")

      pMylq,=plt.plot(N,D_Mylq,label="Loi de Mylq")

    except:

      print "Fit curves seem not to be available"

    plt.legend()

    plt.show()

if __name__=='__main__':

  # Set defaults values

  # Alu can be CPU, GPU or ACCELERATOR

  Alu='CPU'

  # Id of GPU : 1 is for first find !

  Device=0

  # GPU style can be Cuda (Nvidia implementation) or OpenCL

  GpuStyle='OpenCL'

  # Parallel distribution can be on Threads or Blocks

  ParaStyle='Blocks'

  # Iterations is integer

  Iterations=100000000

  # JobStart in first number of Jobs to explore

  JobStart=1

  # JobEnd is last number of Jobs to explore

  JobEnd=16

  # JobStep is the step of Jobs to explore

  JobStep=1

  # Redo is the times to redo the test to improve metrology

  Redo=1

  # OutMetrology is method for duration estimation : False is GPU inside

  OutMetrology=False

  Metrology='InMetro'

  # Curves is True to print the curves

  Curves=False

  # Fit is True to print the curves

  Fit=False

  # DoublePrecision on FP calculus

  DoublePrecision=False

  try:

    opts, args = getopt.getopt(sys.argv[1:],"hoclfa:g:p:i:s:e:t:r:d:",["alu=","gpustyle=","parastyle=","iterations=","jobstart=","jobend=","jobstep=","redo=","device="])

  except getopt.GetoptError:

    print '%s -o (Out of Core Metrology) -c (Print Curves) -l (Double Precision) -f (Fit to Amdahl Law) -a <CPU/GPU/ACCELERATOR> -d <DeviceId> -g <CUDA/OpenCL> -p <Threads/Hybrid/Blocks> -i <Iterations> -s <JobStart> -e <JobEnd> -t <JobStep> -r <RedoToImproveStats> ' % sys.argv[0]

    sys.exit(2)

  for opt, arg in opts:

    if opt == '-h':

      print '%s -o (Out of Core Metrology) -c (Print Curves) -l (Double Precision) -f (Fit to Amdahl Law) -a <CPU/GPU/ACCELERATOR> -d <DeviceId> -g <CUDA/OpenCL> -p <Threads/Hybrid/Blocks> -i <Iterations> -s <JobStart> -e <JobEnd> -t <JobStep> -r <RedoToImproveStats>' % sys.argv[0]

      print "\nInformations about devices detected under OpenCL:"

      # For PyOpenCL import

      try:

        import pyopencl as cl

        Id=1

        for platform in cl.get_platforms():

          for device in platform.get_devices():

            deviceType=cl.device_type.to_string(device.type)

            print "Device #%i of type %s : %s" % (Id,deviceType,device.name)

            Id=Id+1

        print

        sys.exit()

      except ImportError:

        print "Your platform does not seem to support OpenCL"

    elif opt == '-o':

      OutMetrology=True

      Metrology='OutMetro'

    elif opt == '-l':

      DoublePrecision=True

    elif opt == '-c':

      Curves=True

    elif opt == '-f':

      Fit=True

    elif opt in ("-a", "--alu"):

      Alu = arg

    elif opt in ("-d", "--device"):

      Device = int(arg)

    elif opt in ("-g", "--gpustyle"):

      GpuStyle = arg

    elif opt in ("-p", "--parastyle"):

      ParaStyle = arg

    elif opt in ("-i", "--iterations"):

      Iterations = numpy.uint64(arg)

    elif opt in ("-s", "--jobstart"):

      JobStart = int(arg)

    elif opt in ("-e", "--jobend"):

      JobEnd = int(arg)

    elif opt in ("-t", "--jobstep"):

      JobStep = int(arg)

    elif opt in ("-r", "--redo"):

      Redo = int(arg)

  if Alu=='CPU' and GpuStyle=='CUDA':