#!/usr/bin/env python

#
# Splutter-by-MonteCarlo using PyCUDA/PyOpenCL
#
# CC BY-NC-SA 2014 : <emmanuel.quemener@ens-lyon.fr> 
#
# Thanks to Andreas Klockner for PyCUDA:
# http://mathema.tician.de/software/pycuda
# 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

# Marsaglia elements about RNG 

# Common tools
import numpy
from numpy.random import randint as nprnd
import sys
import getopt
import time
import math
from socket import gethostname

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

prout="""

"""


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 CONGfp CONG * 2.328306435454494e-10f
#define SHR3fp SHR3 * 2.328306435454494e-10f
#define MWCfp MWC * 2.328306435454494e-10f
#define KISSfp KISS * 2.328306435454494e-10f

#define MAX (ulong)4294967296
#define UMAX (uint)2147483648

__global__ void SplutterGlobal(uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
    const ulong id=(ulong)(blockIdx.x);
   
    uint z=seed_z-(uint)id;
    uint w=seed_w+(uint)id;

    uint jsr=seed_z;
    uint jcong=seed_w;

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

      // All version
      uint position=(uint)( ((ulong)MWC*(ulong)space)/MAX );

      // UMAX is set to avoid round over overflow
      atomicInc(&s[position],UMAX);
   }

   __syncthreads();
}

__global__ void SplutterGlobalDense(uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
    const ulong id=(ulong)(threadIdx.x+blockIdx.x*blockDim.x);
    const ulong size=(ulong)(gridDim.x*blockDim.x);
    const ulong block=(ulong)space/(ulong)size;
   
    uint z=seed_z-(uint)id;
    uint w=seed_w+(uint)id;

    uint jsr=seed_z;
    uint jcong=seed_w;

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

      // Dense version
       uint position=(uint)( ((ulong)MWC+id*MAX)*block/MAX );

      s[position]++;
   }

   __syncthreads();
}

__global__ void SplutterGlobalSparse(uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{ 
    const ulong id=(ulong)(threadIdx.x+blockIdx.x*blockDim.x);
    const ulong size=(ulong)(gridDim.x*blockDim.x);
    const ulong block=(ulong)space/(ulong)size;
   
    uint z=seed_z-(uint)id;
    uint w=seed_w+(uint)id;

    uint jsr=seed_z;
    uint jcong=seed_w;

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

      // Sparse version
       uint position=(uint)( (ulong)MWC*block/MAX*size+id );

      s[position]++;
   }

   __syncthreads();
}

__global__ void SplutterLocalDense(uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
    const ulong id=(ulong)(threadIdx.x);
    const ulong size=(ulong)(blockDim.x);
    const ulong block=(ulong)space/(ulong)size;
   
    uint z=seed_z-(uint)id;
    uint w=seed_w+(uint)id;

    uint jsr=seed_z;
    uint jcong=seed_w;

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

      // Dense version
       size_t position=(size_t)( ((ulong)MWC+id*MAX)*block/MAX );

      s[position]++;
   }


   __syncthreads();

}

__global__ void SplutterLocalSparse(uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
    const ulong id=(ulong)threadIdx.x;
    const ulong size=(ulong)blockDim.x;
    const ulong block=(ulong)space/(ulong)size;
   
    uint z=seed_z-(uint)id;
    uint w=seed_w+(uint)id;

    uint jsr=seed_z;
    uint jcong=seed_w;

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

      // Sparse version
       size_t position=(size_t)( (ulong)MWC*block/MAX*size+id );

      s[position]++;
   }

   __syncthreads();

}

__global__ void SplutterHybridDense(uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
    const ulong id=(ulong)(blockIdx.x);
    const ulong size=(ulong)(gridDim.x);
    const ulong block=(ulong)space/(ulong)size;
   
    uint z=seed_z-(uint)id;
    uint w=seed_w+(uint)id;

    uint jsr=seed_z;
    uint jcong=seed_w;

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

      // Dense version
      size_t position=(size_t)( ((ulong)MWC+id*MAX)*block/MAX );

      s[position]++;
   }
      
   __syncthreads();
}

__global__ void SplutterHybridSparse(uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
    const ulong id=(ulong)(blockIdx.x);
    const ulong size=(ulong)(gridDim.x);
    const ulong block=(ulong)space/(ulong)size;
   
    uint z=seed_z-(uint)id;
    uint w=seed_w+(uint)id;

    uint jsr=seed_z;
    uint jcong=seed_w;

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

      // Sparse version
      size_t position=(size_t)( (((ulong)MWC*block)/MAX)*size+id );

      s[position]++;

   }

   //s[blockIdx.x]=blockIdx.x;
   __syncthreads();
}

"""

KERNEL_CODE_OPENCL="""
#pragma OPENCL EXTENSION cl_khr_global_int32_base_atomics : 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 CONGfp CONG * 2.328306435454494e-10f
#define SHR3fp SHR3 * 2.328306435454494e-10f
#define MWCfp MWC * 2.328306435454494e-10f
#define KISSfp KISS * 2.328306435454494e-10f

#define MAX (ulong)4294967296

uint rotl(uint value, int shift) {
    return (value << shift) | (value >> (sizeof(value) * CHAR_BIT - shift));
}
 
uint rotr(uint value, int shift) {
    return (value >> shift) | (value << (sizeof(value) * CHAR_BIT - shift));
}

__kernel void SplutterGlobal(__global uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
   __private const ulong id=(ulong)get_global_id(0);
   __private const ulong size=(ulong)get_global_size(0);
   __private const ulong block=(ulong)space/(ulong)size;
   
   __private uint z=seed_z-(uint)id;
   __private uint w=seed_w+(uint)id;

   __private uint jsr=seed_z;
   __private uint jcong=seed_w;

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

      // Dense version
      __private size_t position=(size_t)( ((ulong)MWC*(ulong)space)/MAX );

      atomic_inc(&s[position]);
   }

   barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

}

__kernel void SplutterGlobalDense(__global uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
   __private const ulong id=(ulong)get_global_id(0);
   __private const ulong size=(ulong)get_global_size(0);
   __private const ulong block=(ulong)space/(ulong)size;
   
   __private uint z=seed_z-(uint)id;
   __private uint w=seed_w+(uint)id;

   __private uint jsr=seed_z;
   __private uint jcong=seed_w;

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

      // Dense version
      __private size_t position=(size_t)( ((ulong)MWC+id*MAX)*block/MAX );

      s[position]++;
   }

   barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

}

__kernel void SplutterGlobalSparse(__global uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
   __private const ulong id=(ulong)get_global_id(0);
   __private const ulong size=(ulong)get_global_size(0);
   __private const ulong block=(ulong)space/(ulong)size;
   
   __private uint z=seed_z-(uint)id;
   __private uint w=seed_w+(uint)id;

   __private uint jsr=seed_z;
   __private uint jcong=seed_w;

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

      // Sparse version
      __private size_t position=(size_t)( (ulong)MWC*block/MAX*size+id );

      s[position]++;
   }

   barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

}

__kernel void SplutterLocalDense(__global uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
   __private const ulong id=(ulong)get_local_id(0);
   __private const ulong size=(ulong)get_local_size(0);
   __private const ulong block=(ulong)space/(ulong)size;
   
   __private uint z=seed_z-(uint)id;
   __private uint w=seed_w+(uint)id;

   __private uint jsr=seed_z;
   __private uint jcong=seed_w;

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

      // Dense version
      __private size_t position=(size_t)( ((ulong)MWC+id*MAX)*block/MAX );

      s[position]++;
   }

   barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

}

__kernel void SplutterLocalSparse(__global uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
   __private const ulong id=(ulong)get_local_id(0);
   __private const ulong size=(ulong)get_local_size(0);
   __private const ulong block=(ulong)space/(ulong)size;
   
   __private uint z=seed_z-(uint)id;
   __private uint w=seed_w+(uint)id;

   __private uint jsr=seed_z;
   __private uint jcong=seed_w;

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

      // Sparse version
      __private size_t position=(size_t)( (ulong)MWC*block/MAX*size+id );

      s[position]++;
   }

   barrier(CLK_LOCAL_MEM_FENCE | CLK_GLOBAL_MEM_FENCE);

}

__kernel void SplutterHybridDense(__global uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
   __private const ulong id=(ulong)(get_global_id(0));
   __private const ulong size=(ulong)(get_local_size(0)*get_num_groups(0));
   __private const ulong block=(ulong)space/(ulong)size;
   
   __private uint z=seed_z-(uint)id;
   __private uint w=seed_w+(uint)id;

   __private uint jsr=seed_z;
   __private uint jcong=seed_w;

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

      // Dense version
      __private size_t position=(size_t)( ((ulong)MWC+id*MAX)*block/MAX );

      s[position]++;
   }
      
}

__kernel void SplutterHybridSparse(__global uint *s,const uint space,const ulong iterations,const uint seed_w,const uint seed_z)
{
   __private const ulong id=(ulong)(get_global_id(0));
   __private const ulong size=(ulong)(get_local_size(0)*get_num_groups(0));
   __private const ulong block=(ulong)space/(ulong)size;
   
   __private uint z=seed_z-(uint)id;
   __private uint w=seed_w+(uint)id;

   __private uint jsr=seed_z;
   __private uint jcong=seed_w;

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

      // Sparse version
      __private size_t position=(size_t)( (ulong)MWC*block/MAX*size+id );

      s[position]++;
   }
      
}

"""

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

  # Avec PyCUDA autoinit, rien a faire !

  circleCU = cuda.InOut(circle)
  
  mod = SourceModule(KERNEL_CODE_CUDA)

  if Density=='Dense':
    MetropolisBlocksCU=mod.get_function("SplutterGlobalDense")
    MetropolisThreadsCU=mod.get_function("SplutterLocalDense")
    MetropolisHybridCU=mod.get_function("SplutterHybridDense")
  elif Density=='Sparse':
    MetropolisBlocksCU=mod.get_function("SplutterGlobalSparse")
    MetropolisThreadsCU=mod.get_function("SplutterLocalSparse")
    MetropolisHybridCU=mod.get_function("SplutterHybridSparse")
  else:
    MetropolisBlocksCU=mod.get_function("SplutterGlobal")
    
  start = pycuda.driver.Event()
  stop = pycuda.driver.Event()
  
  MySplutter=numpy.zeros(steps)
  MyDuration=numpy.zeros(steps)

  if iterations%jobs==0:
    iterationsCL=numpy.uint64(iterations/jobs)
  else:
    iterationsCL=numpy.uint64(iterations/jobs+1)
    
  iterationsNew=iterationsCL*jobs

  Splutter=numpy.zeros(jobs*16).astype(numpy.uint32)

  for i in range(steps):

    Splutter[:]=0
    
    print Splutter,len(Splutter)

    SplutterCU = cuda.InOut(Splutter)

    start.record()
    start.synchronize()
    if ParaStyle=='Blocks':
      MetropolisBlocksCU(SplutterCU,
                         numpy.uint32(len(Splutter)),
                         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)
      MetropolisHybridCU(SplutterCU,
                         numpy.uint32(len(Splutter)),
                         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:
      MetropolisThreadsCU(SplutterCU,
                       numpy.uint32(len(Splutter)),
                       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,1,jobs,ParaStyle)
    stop.record()
    stop.synchronize()
                
    elapsed = start.time_till(stop)*1e-3

    print Splutter,sum(Splutter)
    MySplutter[i]=numpy.median(Splutter)
    print numpy.mean(Splutter),MySplutter[i],numpy.std(Splutter)

    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,Density):
	
  # Initialisation des variables en les CASTant correctement

  MaxMemoryXPU=0
  MinMemoryXPU=0

  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)
        deviceMemory=device.max_mem_alloc_size
        if deviceMemory>MaxMemoryXPU:
          MaxMemoryXPU=deviceMemory
        if deviceMemory<MinMemoryXPU or MinMemoryXPU==0:
          MinMemoryXPU=deviceMemory
        if deviceType=="GPU" and Alu=="GPU" and not HasXPU:
          XPU=device
          print "GPU selected with Allocable Memory %i: %s" % (deviceMemory,device.name)
          HasXPU=True
          MemoryXPU=deviceMemory
        if deviceType=="CPU" and Alu=="CPU" and not HasXPU:        
          XPU=device
          print "CPU selected with Allocable Memory %i: %s" % (deviceMemory,device.name)
          HasXPU=True
          MemoryXPU=deviceMemory
          
  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)
        deviceMemory=device.max_mem_alloc_size
        if deviceMemory>MaxMemoryXPU:
          MaxMemoryXPU=deviceMemory
        if deviceMemory<MinMemoryXPU or MinMemoryXPU==0:
          MinMemoryXPU=deviceMemory
        if Id==Device and Alu==deviceType and HasXPU==False:
          XPU=device
          print "CPU/GPU selected with Allocable Memory %i: %s" % (deviceMemory,device.name)
          HasXPU=True
          MemoryXPU=deviceMemory
        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)

  print "Allocable Memory is %i, between %i and %i " % (MemoryXPU,MinMemoryXPU,MaxMemoryXPU)

  # 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

  MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build(options = "-cl-mad-enable -cl-fast-relaxed-math")
      
  MyDuration=numpy.zeros(steps)
  
  if iterations%jobs==0:
    iterationsCL=numpy.uint64(iterations/jobs)
  else:
    iterationsCL=numpy.uint64(iterations/jobs+1)
    
  iterationsNew=numpy.uint64(iterationsCL*jobs)

  MySplutter=numpy.zeros(steps)

  MaxWorks=2**(int)(numpy.log2(MinMemoryXPU/4))
  print MaxWorks,2**(int)(numpy.log2(MemoryXPU))
  
  #Splutter=numpy.zeros((MaxWorks/jobs)*jobs).astype(numpy.uint32)
  Splutter=numpy.zeros(jobs*16).astype(numpy.uint32)

  for i in range(steps):
		
    #Splutter=numpy.zeros(2**(int)(numpy.log2(MemoryXPU/4))).astype(numpy.uint32)
    #Splutter=numpy.zeros(1024).astype(numpy.uint32)
 
    #Splutter=numpy.zeros(jobs).astype(numpy.uint32)

    Splutter[:]=0

    print Splutter,len(Splutter)

    SplutterCL = cl.Buffer(ctx, mf.WRITE_ONLY|mf.COPY_HOST_PTR,hostbuf=Splutter)

    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
      # CLLaunch=MetropolisCL.MainLoopGlobal(queue,(jobs,),None,
      #                                      SplutterCL,
      #                                      numpy.uint32(len(Splutter)),
      #                                      numpy.uint64(iterationsCL),
      #                                      numpy.uint32(nprnd(2**30/jobs)),
      #                                      numpy.uint32(nprnd(2**30/jobs)))
      if Density=='Dense':
        CLLaunch=MetropolisCL.SplutterGlobalDense(queue,(jobs,),None,
                                                  SplutterCL,
                                                  numpy.uint32(len(Splutter)),
                                                  numpy.uint64(iterationsCL),
                                                  numpy.uint32(521288629),
                                                  numpy.uint32(362436069))
      elif Density=='Sparse':
        CLLaunch=MetropolisCL.SplutterGlobalSparse(queue,(jobs,),None,
                                                   SplutterCL,
                                                   numpy.uint32(len(Splutter)),
                                                   numpy.uint64(iterationsCL),
                                                   numpy.uint32(521288629),
                                                   numpy.uint32(362436069))
        
      else:
        CLLaunch=MetropolisCL.SplutterGlobal(queue,(jobs,),None,
                                             SplutterCL,
                                             numpy.uint32(len(Splutter)),
                                             numpy.uint64(iterationsCL),
                                             numpy.uint32(521288629),
                                             numpy.uint32(362436069))
        
      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 Density=='Dense':
        CLLaunch=MetropolisCL.SplutterHybridDense(queue,(jobs,),(threads,),
                                                  SplutterCL,
                                                  numpy.uint32(len(Splutter)),
                                                  numpy.uint64(iterationsCL),
                                                  numpy.uint32(nprnd(2**30/jobs)),
                                                  numpy.uint32(nprnd(2**30/jobs)))
      elif Density=='Sparse':
        CLLaunch=MetropolisCL.SplutterHybridSparse(queue,(jobs,),(threads,),
                                                   SplutterCL,
                                                   numpy.uint32(len(Splutter)),
                                                   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 Density=='Dense':
        CLLaunch=MetropolisCL.SplutterLocalDense(queue,(jobs,),(jobs,),
                                                 SplutterCL,
                                                 numpy.uint32(len(Splutter)),
                                                 numpy.uint64(iterationsCL),
                                                 numpy.uint32(nprnd(2**30/jobs)),
                                                 numpy.uint32(nprnd(2**30/jobs)))
      elif Density=='Sparse':
        CLLaunch=MetropolisCL.SplutterLocalSparse(queue,(jobs,),(jobs,),
                                                  SplutterCL,
                                                  numpy.uint32(len(Splutter)),
                                                  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, Splutter, SplutterCL).wait()

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

    MyDuration[i]=elapsed
    print Splutter,sum(Splutter)
    #MySplutter[i]=numpy.median(Splutter)
    #print numpy.mean(Splutter)*len(Splutter),MySplutter[i]*len(Splutter),numpy.std(Splutter)

    SplutterCL.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=10000000
  # 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
  # Spluttering is Dense by default
  Density='All'

  try:
    opts, args = getopt.getopt(sys.argv[1:],"hocfa:g:p:i:s:e:t:r:d:y:",["alu=","gpustyle=","parastyle=","iterations=","jobstart=","jobend=","jobstep=","redo=","device=","density="])
  except getopt.GetoptError:
    print '%s -o (Out of Core Metrology) -c (Print Curves) -f (Fit to Amdahl Law) -y <Dense/Sparse/All> -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) -f (Fit to Amdahl Law)  -y <Dense/Sparse/All> -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)
            deviceMemory=device.max_mem_alloc_size
            print "Device #%i of type %s with memory %i : %s" % (Id,deviceType,deviceMemory,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 == '-c':
      Curves=True
    elif opt in ("-y", "--density"):
      Density = arg
    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)

  print "Toto %s" % Alu

  if Alu=='CPU' and GpuStyle=='CUDA':
    print "Alu can't be CPU for CUDA, set Alu to GPU"
    Alu='GPU'

  if ParaStyle not in ('Blocks','Threads','Hybrid'):
    print "%s not exists, ParaStyle set as Threads !" % ParaStyle
    ParaStyle='Blocks'

  print "Compute unit : %s" % Alu
  print "Device Identification : %s" % Device
  print "GpuStyle used : %s" % GpuStyle
  print "Parallel Style used : %s" % ParaStyle
  print "Density Spluttering : %s" % Density
  print "Iterations : %s" % Iterations
  print "Number of threads on start : %s" % JobStart
  print "Number of threads on end : %s" % JobEnd
  print "Number of redo : %s" % Redo
  print "Metrology done out of CPU/GPU : %r" % OutMetrology

  if GpuStyle=='CUDA':
    try:
      # For PyCUDA import
      import pycuda.driver as cuda
      import pycuda.gpuarray as gpuarray
      import pycuda.autoinit
      from pycuda.compiler import SourceModule
    except ImportError:
      print "Platform does not seem to support CUDA"

  if GpuStyle=='OpenCL':
    try:
      # For PyOpenCL import
      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)
          if Id == Device:
            # Set the Alu as detected Device Type
            Alu=deviceType
          Id=Id+1
    except ImportError:
      print "Platform does not seem to support CUDA"
      
  average=numpy.array([]).astype(numpy.float32)
  median=numpy.array([]).astype(numpy.float32)
  stddev=numpy.array([]).astype(numpy.float32)

  ExploredJobs=numpy.array([]).astype(numpy.uint32)

  Jobs=JobStart

  while Jobs <= JobEnd:
    avg,med,std=0,0,0
    ExploredJobs=numpy.append(ExploredJobs,Jobs)
    circle=numpy.zeros(Jobs).astype(numpy.uint64)

    if OutMetrology:
      duration=numpy.array([]).astype(numpy.float32)
      for i in range(Redo):
        start=time.time()
        if GpuStyle=='CUDA':
          try:
            a,m,s=MetropolisCuda(circle,Iterations,1,Jobs,ParaStyle,Density)
          except:
            print "Problem with %i // computations on Cuda" % Jobs
        elif GpuStyle=='OpenCL':
          try:
            a,m,s=MetropolisOpenCL(circle,Iterations,1,Jobs,ParaStyle,
                                   Alu,Device,Density)
          except:
            print "Problem with %i // computations on OpenCL" % Jobs            
        duration=numpy.append(duration,time.time()-start)
      if (a,m,s) != (0,0,0):
        avg=numpy.mean(duration)
        med=numpy.median(duration)
        std=numpy.std(duration)
      else:
        print "Values seem to be wrong..."
    else:
      if GpuStyle=='CUDA':
        try:
          avg,med,std=MetropolisCuda(circle,Iterations,Redo,
                                     Jobs,ParaStyle,Density)
        except:
          print "Problem with %i // computations on Cuda" % Jobs
      elif GpuStyle=='OpenCL':
        try:
          avg,med,std=MetropolisOpenCL(circle,Iterations,Redo,
                                       Jobs,ParaStyle,Alu,Device,Density)
        except:
          print "Problem with %i // computations on OpenCL" % Jobs            

    if (avg,med,std) != (0,0,0):
      print "jobs,avg,med,std",Jobs,avg,med,std
      average=numpy.append(average,avg)
      median=numpy.append(median,med)
      stddev=numpy.append(stddev,std)
    else:
      print "Values seem to be wrong..."
    #THREADS*=2
    if len(average)!=0:
      numpy.savez("Splutter_%s_%s_%s_%i_%i_%.8i_Device%i_%s_%s" % (Alu,GpuStyle,ParaStyle,JobStart,JobEnd,Iterations,Device,Metrology,gethostname()),(ExploredJobs,average,median,stddev))
      ToSave=[ ExploredJobs,average,median,stddev ]
      numpy.savetxt("Splutter_%s_%s_%s_%i_%i_%.8i_Device%i_%s_%s" % (Alu,GpuStyle,ParaStyle,JobStart,JobEnd,Iterations,Device,Metrology,gethostname()),numpy.transpose(ToSave))
    Jobs+=JobStep

  if Fit:
    FitAndPrint(ExploredJobs,median,Curves)