Index: BLAS/xGEMM/fortran_common.h.orig
===================================================================
--- BLAS/xGEMM/fortran_common.h.orig	(revision 6)
+++ BLAS/xGEMM/fortran_common.h.orig	(revision 7)
@@ -1,5 +1,5 @@
 /*
- * Copyright 1993-2011 NVIDIA Corporation.  All rights reserved.
+ * Copyright 1993-2012 NVIDIA Corporation.  All rights reserved.
  *
  * NOTICE TO LICENSEE:
  *
@@ -46,24 +46,30 @@
  * comments to the code, the above Disclaimer and U.S. Government End
  * Users Notice.
  */
- 
-#define CUBLAS_G77              1
-#define CUBLAS_INTEL_FORTRAN    2
-#define CUBLAS_G95              3
 
-/* Default to g77 on Linux, and Intel Fortran on Win32 */
+
+
+/* Define CUBLAS_FORTRAN_COMPILER for Windows needed because 
+   tbe compilation is done from the perl script fortran_nightly.pl which
+   does not include fortran_tools.mk
+   On Linux and Apple, CFLAGS is setup with fortran_tools.mk
+*/
 #if defined(_WIN32)
-#define CUBLAS_FORTRAN_COMPILER CUBLAS_INTEL_FORTRAN
-#elif defined(__linux)
-#define CUBLAS_FORTRAN_COMPILER CUBLAS_G95
-#elif defined(__APPLE__)
-#define CUBLAS_FORTRAN_COMPILER CUBLAS_G95
-#define RETURN_COMPLEX   1
-#else
-#error unsupported platform
+#ifndef CUBLAS_INTEL_FORTRAN
+#define CUBLAS_INTEL_FORTRAN
 #endif
+#endif
 
-#if (CUBLAS_FORTRAN_COMPILER==CUBLAS_G77) || (CUBLAS_FORTRAN_COMPILER==CUBLAS_G95)
+#if defined(CUBLAS_GFORTRAN)
+/* using option -ff2c make the ABI compatible with F77
+   No need to define RETURN_COMPLEX, which cause problem 
+   on Gfortran 4.x on 32 bit
+*/
+/* #define RETURN_COMPLEX 1 */
+#endif
+
+
+#if defined(CUBLAS_GFORTRAN) || defined (CUBLAS_G95)
 /* NOTE: Must use -fno-second-underscore when building Fortran source with g77
  *       g77 invocation may not use -fno-f2c, which forces different return 
  *       type conventions than the one used below
@@ -231,7 +237,7 @@
 #define CUBLAS_ZHER2           cublas_zher2_
 #define CUBLAS_ZHPR2           cublas_zhpr2_
 
-#elif CUBLAS_FORTRAN_COMPILER==CUBLAS_INTEL_FORTRAN
+#elif defined(CUBLAS_INTEL_FORTRAN)
 
 #define CUBLAS_INIT             CUBLAS_INIT 
 #define CUBLAS_SHUTDOWN         CUBLAS_SHUTDOWN
Index: BLAS/xGEMM/Makefile
===================================================================
--- BLAS/xGEMM/Makefile	(revision 6)
+++ BLAS/xGEMM/Makefile	(revision 7)
@@ -1,6 +1,7 @@
 SOURCE=xGEMM.c
 
-CC=gcc
+CC=gcc-4.6
+FC=gfortran-4.6
 CFLAGS=-Wall -O3
 LDFLAGS=-lm
 CUDADIR=/opt/cuda
@@ -16,14 +17,15 @@
 
 GOTO2=/opt/GotoBLAS2
 
-OPENBLAS=/opt/OpenBLAS
+OPENBLAS=/usr/lib/openblas-base
 
 ACML=/opt/acml
 ACMLINC=$(ACML)/gfortran64_mp/include
 ACMLLIB=$(ACML)/gfortran64_mp/lib
 
+ATLAS=/usr/lib/atlas-base
 
-EXECUTABLE=cblas fblas gsl cublas thunking gotoblas openblas acml
+EXECUTABLE=cblas fblas gsl cublas thunking openblas
 
 FORMAT=DOUBLE
 #FORMAT=FLOAT
@@ -36,11 +38,11 @@
 
 cblas: $(SOURCE)
 
-	$(CC) $(CFLAGS) $(DIRECTIVES) -DFLOAT -DCBLAS $(LDFLAGS) \
-		$(SOURCE) -lcblas -o $(SOURCE:.c=)_SP_$@
+	$(CC) -L$(ATLAS) $(CFLAGS) $(DIRECTIVES) -DFLOAT -DCBLAS $(LDFLAGS) \
+		$(SOURCE) -latlas -lcblas -o $(SOURCE:.c=)_SP_$@
 
-	$(CC) $(CFLAGS) $(DIRECTIVES) -DDOUBLE -DCBLAS $(LDFLAGS) \
-		$(SOURCE) -lcblas -o $(SOURCE:.c=)_DP_$@
+	$(CC) -L$(ATLAS) $(CFLAGS) $(DIRECTIVES) -DDOUBLE -DCBLAS $(LDFLAGS) \
+		$(SOURCE) -latlas -lcblas -o $(SOURCE:.c=)_DP_$@
 
 gotoblas: $(SOURCE)
 
@@ -70,11 +72,11 @@
 
 fblas: $(SOURCE)
 
-	$(CC) $(CFLAGS) $(DIRECTIVES) -DFLOAT -DFBLAS $(LDFLAGS) \
-		$(SOURCE) -lf77blas -o $(SOURCE:.c=)_SP_$@
+	$(CC) -L$(ATLAS) $(CFLAGS) $(DIRECTIVES) -DFLOAT -DFBLAS $(LDFLAGS) \
+		$(SOURCE) -latlas -lf77blas -o $(SOURCE:.c=)_SP_$@
 
-	$(CC) $(CFLAGS) $(DIRECTIVES) -DDOUBLE -DFBLAS $(LDFLAGS) \
-		$(SOURCE) -lf77blas -o $(SOURCE:.c=)_DP_$@
+	$(CC) -L$(ATLAS) $(CFLAGS) $(DIRECTIVES) -DDOUBLE -DFBLAS $(LDFLAGS) \
+		$(SOURCE) -latlas -lf77blas -o $(SOURCE:.c=)_DP_$@
 
 
 gsl: $(SOURCE)
@@ -103,14 +105,15 @@
 	cp $(CUDASRC)/$(THUNKING) $(CUDASRC)/$(THUNKING:.c=.h) \
 		$(CUDASRC)/$(CUDASRCINC) .
 	# Patch Thunking prototypes to compile on Debian Lenny
-	patch $(CUDASRCINC) $(PATCHTHUNKING)
-	$(CC) -I$(CUDAINC) $(CFLAGS) -c $(THUNKING)
+	$(CC) -DCUBLAS_GFORTRAN -I$(CUDAINC) $(CFLAGS) -c $(THUNKING)
 
-	$(CC) -I$(CUDAINC) -L$(CUDALIB) $(CFLAGS) -DFLOAT -DTHUNKING \
+	$(CC) -DCUBLAS_GFORTRAN \
+		-I$(CUDAINC) -L$(CUDALIB) $(CFLAGS) -DFLOAT -DTHUNKING \
 		$(LDFLAGS) $(DIRECTIVES) $(SOURCE) \
 		$(THUNKING:.c=.o) -lcublas -o $(SOURCE:.c=)_SP_$@
 
-	$(CC) -I$(CUDAINC) -L$(CUDALIB) $(CFLAGS) -DDOUBLE -DTHUNKING \
+	$(CC) -DCUBLAS_GFORTRAN \
+		-I$(CUDAINC) -L$(CUDALIB) $(CFLAGS) -DDOUBLE -DTHUNKING \
 		$(LDFLAGS) $(DIRECTIVES) $(SOURCE) \
 		$(THUNKING:.c=.o) -lcublas -o $(SOURCE:.c=)_DP_$@
 
Index: BLAS/xGEMM/xGEMM.c
===================================================================
--- BLAS/xGEMM/xGEMM.c	(revision 6)
+++ BLAS/xGEMM/xGEMM.c	(revision 7)
@@ -19,6 +19,8 @@
 #define CUBLAS_WRAPPER_ERROR_STUB       4
 #elif THUNKING
 #include <cublas.h>
+#include "fortran_common.h"
+#include "fortran_thunking.h"
 #elif FBLAS
 #include <cblas_f77.h>
 #elif GSL
@@ -36,11 +38,8 @@
 #define LENGTH float
 #endif
 
-#ifdef THUNKING
-#include "fortran_thunking.h"
+#ifdef FBLAS
 
-#elif FBLAS
-
 #ifdef DOUBLE
 
 void dgemm_(FCHAR, FCHAR, FINT, FINT, FINT, const double *, const double *, FINT, 
@@ -52,7 +51,6 @@
 	       const float *, FINT, const float *, float *, FINT);
 
 #endif
-
 #endif
 
 /* Matrix with only defined triangular terms */
@@ -320,9 +318,9 @@
     {      
       /* Multiply A by X as Y <- A.X */
       CUBLAS_DGEMM(&transa,&transa,
-		   &dim,&dim,&dim,&alpha,B,&dim,A,&dim,&beta,C,&dim);
+      		   &dim,&dim,&dim,&alpha,B,&dim,A,&dim,&beta,C,&dim);
       CUBLAS_DGEMM(&transb,&transb,
-		   &dim,&dim,&dim,&alpha,A,&dim,B,&dim,&beta,D,&dim);
+      		   &dim,&dim,&dim,&alpha,A,&dim,B,&dim,&beta,D,&dim);
     }
 
 #else
@@ -331,9 +329,9 @@
     {      
       /* Multiply A by X as Y <- A.X */
       CUBLAS_SGEMM(&transa,&transa,
-		   &dim,&dim,&dim,&alpha,B,&dim,A,&dim,&beta,C,&dim);
+      		   &dim,&dim,&dim,&alpha,B,&dim,A,&dim,&beta,C,&dim);
       CUBLAS_SGEMM(&transb,&transb,
-		   &dim,&dim,&dim,&alpha,A,&dim,B,&dim,&beta,D,&dim);
+      		   &dim,&dim,&dim,&alpha,A,&dim,B,&dim,&beta,D,&dim);
     }
   
 #endif
Index: BLAS/xTRSV/Makefile
===================================================================
--- BLAS/xTRSV/Makefile	(revision 6)
+++ BLAS/xTRSV/Makefile	(revision 7)
@@ -1,6 +1,7 @@
 SOURCE=xTRSV.c
 
-CC=gcc
+CC=gcc-4.6
+FC=gfortran-4.6
 CFLAGS=-Wall -O3
 LDFLAGS=-lm
 CUDADIR=/opt/cuda
@@ -10,19 +11,20 @@
 CUDAINC=$(CUDADIR)/include
 CUDALIB=$(CUDADIR)/lib64
 
-PATCHTHUNKING=patch_thunking.h
-
 GSLINC=/usr/include/gsl
 
 GOTO2=/opt/GotoBLAS2
 
-OPENBLAS=/opt/OpenBLAS
+OPENBLAS=/usr/lib/openblas-base
 
+ATLAS=/usr/lib/atlas-base
+
 ACML=/opt/acml
 ACMLINC=$(ACML)/gfortran64_mp/include
 ACMLLIB=$(ACML)/gfortran64_mp/lib
 
-EXECUTABLE=cblas fblas gsl cublas thunking gotoblas acml openblas
+#EXECUTABLE=cblas fblas gsl cublas thunking gotoblas acml openblas
+EXECUTABLE=cblas fblas gsl cublas thunking openblas
 
 #FORMAT=DOUBLE
 FORMAT=FLOAT
@@ -36,10 +38,10 @@
 cblas: $(SOURCE)
 
 	$(CC) $(CFLAGS) $(DIRECTIVES) -DFLOAT -DCBLAS $(LDFLAGS) \
-		$(SOURCE) -lcblas -o $(SOURCE:.c=)_SP_$@
+		$(SOURCE) $(ATLAS)/libcblas.a $(ATLAS)/libatlas.a -o $(SOURCE:.c=)_SP_$@
 
 	$(CC) $(CFLAGS) $(DIRECTIVES) -DDOUBLE -DCBLAS $(LDFLAGS) \
-		$(SOURCE) -lcblas -o $(SOURCE:.c=)_DP_$@
+		$(SOURCE) $(ATLAS)/libcblas.a $(ATLAS)/libatlas.a -o $(SOURCE:.c=)_DP_$@
 
 gotoblas: $(SOURCE)
 
@@ -70,10 +72,10 @@
 fblas: $(SOURCE)
 
 	$(CC) $(CFLAGS) $(DIRECTIVES) -DFLOAT -DFBLAS $(LDFLAGS) \
-		$(SOURCE) -lf77blas -o $(SOURCE:.c=)_SP_$@
+		$(SOURCE) $(ATLAS)/libf77blas.a $(ATLAS)/libatlas.a -lgfortran -o $(SOURCE:.c=)_SP_$@
 
 	$(CC) $(CFLAGS) $(DIRECTIVES) -DDOUBLE -DFBLAS $(LDFLAGS) \
-		$(SOURCE) -lf77blas -o $(SOURCE:.c=)_DP_$@
+		$(SOURCE) $(ATLAS)/libf77blas.a $(ATLAS)/libatlas.a -lgfortran -o $(SOURCE:.c=)_DP_$@
 
 
 gsl: $(SOURCE)
@@ -98,15 +100,14 @@
 # Copy of source for thunking CUBLAS approach
 	cp $(CUDASRC)/$(THUNKING) $(CUDASRC)/$(THUNKING:.c=.h) \
 		$(CUDASRC)/$(CUDASRCINC) .
-# Patch Thunking prototypes to compile on Debian Lenny
-	patch $(CUDASRCINC) $(PATCHTHUNKING)
-	$(CC) -I$(CUDAINC) $(CFLAGS) -c $(THUNKING)
 
-	$(CC) -I$(CUDAINC) -L$(CUDALIB) $(CFLAGS) -DTHUNKING \
+	$(CC) -DCUBLAS_GFORTRAN -I$(CUDAINC) $(CFLAGS) -c $(THUNKING)
+
+	$(CC) -DCUBLAS_GFORTRAN -I$(CUDAINC) -L$(CUDALIB) $(CFLAGS) -DTHUNKING \
 		$(LDFLAGS) $(DIRECTIVES) $(SOURCE) -DFLOAT \
 		$(THUNKING:.c=.o) -lcublas -o $(SOURCE:.c=)_SP_$@
 
-	$(CC) -I$(CUDAINC) -L$(CUDALIB) $(CFLAGS) -DTHUNKING \
+	$(CC) -DCUBLAS_GFORTRAN -I$(CUDAINC) -L$(CUDALIB) $(CFLAGS) -DTHUNKING \
 		$(LDFLAGS) $(DIRECTIVES) $(SOURCE) -DDOUBLE \
 		$(THUNKING:.c=.o) -lcublas -o $(SOURCE:.c=)_DP_$@
 
Index: BLAS/xTRSV/xTRSV.c
===================================================================
--- BLAS/xTRSV/xTRSV.c	(revision 6)
+++ BLAS/xTRSV/xTRSV.c	(revision 7)
@@ -22,6 +22,8 @@
 #define CUBLAS_WRAPPER_ERROR_STUB       4
 #elif THUNKING
 #include <cublas.h>
+#include "fortran_common.h"
+#include "fortran_thunking.h"
 #elif FBLAS
 #include <cblas_f77.h>
 #elif GSL
@@ -39,63 +41,10 @@
 #define LENGTH float
 #endif
 
-#ifdef THUNKING
-/* WARNING !
-Prototypes from fortran.c functions used MUST be defined here !
-*/
-#include "fortran_thunking.h"
+#ifdef FBLAS
 
-/*
 #ifdef DOUBLE
 
-void CUBLAS_DCOPY (const int *n, const double *x, const int *incx, double *y,
-                   const int *incy);
-
-double CUBLAS_DNRM2 (const int *dim, const double *X, const int *incx);
-
-void CUBLAS_DTRSV (const char *uplo, const char *trans, const char *diag,
-                   const int *n, const double *A, const int *lda, double *x,
-                   const int *incx);
-
-void CUBLAS_DGEMV (const char *trans, const int *m, const int *n,
-                   const double *alpha, const double *A, const int *lda,
-                   const double *x, const int *incx, const double *beta,
-                   double *y, const int *incy);
-
-void CUBLAS_DSWAP (const int *n, double *x, const int *incx, double *y,
-                   const int *incy);
-
-void CUBLAS_DAXPY (const int *n, const double *alpha, const double *x, 
-                   const int *incx, double *y, const int *incy);
-
-#else
-void CUBLAS_SCOPY (const int *n, const float *x, const int *incx, float *y,
-                   const int *incy);
-
-float CUBLAS_SNRM2 (const int *dim, const float *X, const int *incx);
-
-void CUBLAS_STRSV (const char *uplo, const char *trans, const char *diag,
-                   const int *n, const float *A, const int *lda, float *x,
-                   const int *incx);
-
-void CUBLAS_SGEMV (const char *trans, const int *m, const int *n,
-                   const float *alpha, const float *A, const int *lda,
-                   const float *x, const int *incx, const float *beta,
-                   float *y, const int *incy);
-
-void CUBLAS_SSWAP (const int *n, float *x, const int *incx, float *y,
-                   const int *incy);
-
-void CUBLAS_SAXPY (const int *n, const float *alpha, const float *x, 
-                   const int *incx, float *y, const int *incy);
-
-#endif
-*/
-
-#elif FBLAS
-
-#ifdef DOUBLE
-
 void dtrsv_( FCHAR, FCHAR, FCHAR, FINT, const double *, FINT, double *, FINT);
 
 void dgemv_(FCHAR, FINT, FINT, const double *, const double *, FINT, 
Index: Pi/GPU/Pi-GPU.py-20130213-2
===================================================================
--- Pi/GPU/Pi-GPU.py-20130213-2	(revision 0)
+++ Pi/GPU/Pi-GPU.py-20130213-2	(revision 7)
@@ -0,0 +1,602 @@
+#!/usr/bin/env python
+
+#
+# Pi-by-MC 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 matplotlib.pyplot as plt
+import math
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+  return (T1*(s+c1*N+c2*N*N+p/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(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(blockIdx.x+1);
+   uint w=seed_w/(blockIdx.x+1);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopThreads(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(threadIdx.x+1);
+   uint w=seed_w/(threadIdx.x+1);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[threadIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopHybrid(uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockDim.x*blockIdx.x+threadIdx.x]=total;
+   __syncthreads();
+
+}
+"""
+
+KERNEL_CODE_OPENCL="""
+
+// 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 uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+   
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+     float y=MWCfp ;
+
+      // Matching test
+      int 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):
+
+  # 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")
+  
+  start = pycuda.driver.Event()
+  stop = pycuda.driver.Event()
+  
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%jobs==0:
+    iterationsCL=numpy.uint32(iterations/jobs+1)
+    iterationsNew=iterationsCL*jobs
+  else:
+    iterationsCL=numpy.uint32(iterations/jobs)
+    iterationsNew=iterations
+
+  for i in range(steps):
+    start.record()
+    start.synchronize()
+    if ParaStyle=='Blocks':
+      MetropolisBlocksCU(circleCU,
+                         numpy.uint32(iterationsCL),
+                         numpy.uint32(nprnd(2**32/jobs)),
+                         numpy.uint32(nprnd(2**32/jobs)),
+                         grid=(jobs,1),
+                         block=(1,1,1))
+      print "GPU with %i %s done" % (jobs,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      blocks=jobs/int(math.sqrt(float(jobs)))
+      MetropolisHybridCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          grid=(blocks,1),
+                          block=(jobs/blocks,1,1))
+      print "GPU with (blocks,jobs)=(%i,%i) %s done" % (blocks,jobs/blocks,ParaStyle)
+    else:
+      MetropolisJobsCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          grid=(1,1),
+                          block=(jobs,1,1))
+      print "GPU with %i %s done" % (jobs,ParaStyle)
+    stop.record()
+    stop.synchronize()
+                
+    #elapsed = stop.time_since(start)*1e-3
+    elapsed = start.time_till(stop)*1e-3
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsCL)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+    #time.sleep(1)
+
+  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):
+	
+  # Initialisation des variables en les CASTant correctement
+    
+  # Je detecte un peripherique GPU dans la liste des peripheriques
+  # for platform in cl.get_platforms():
+  # 	for device in platform.get_devices():
+  # 		if cl.device_type.to_string(device.type)=='GPU':
+  # 			GPU=device
+  #print "GPU detected: ",device.name
+
+  HasGPU=False
+  Id=1
+  # Device selection based on choice (default is GPU)
+  for platform in cl.get_platforms():
+    for device in platform.get_devices():
+      if not HasGPU:
+        deviceType=cl.device_type.to_string(device.type)
+        if deviceType=="GPU" and Alu=="GPU" and Id==Device:
+          GPU=device
+          print "GPU selected: ",device.name
+          HasGPU=True
+        if deviceType=="CPU" and Alu=="CPU":	    
+          GPU=device
+          print "CPU selected: ",device.name
+          HasGPU=True
+      Id=Id+1
+				
+  # Je cree le contexte et la queue pour son execution
+  #ctx = cl.create_some_context()
+  ctx = cl.Context([GPU])
+  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")
+
+  #MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build()
+
+  i=0
+
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%jobs==0:
+    iterationsCL=numpy.uint32(iterations/jobs+1)
+    iterationsNew=iterationsCL*jobs
+  else:
+    iterationsCL=numpy.uint32(iterations/jobs)
+    iterationsNew=iterations
+
+  blocks=int(math.sqrt(jobs))
+
+  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
+      CLLaunch=MetropolisCL.MainLoopGlobal(queue,(jobs,),None,
+                                           circleCL,
+                                           numpy.uint32(iterationsCL),
+                                           numpy.uint32(nprnd(2**32/jobs)),
+                                           numpy.uint32(nprnd(2**32/jobs)))
+      print "%s with %i %s done" % (Alu,jobs,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopHybrid(queue,(blocks*blocks,),(blocks,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/jobs)),
+                                          numpy.uint32(nprnd(2**32/jobs)))
+      print "%s with (Blocks,Threads)=(%i,%i) %s done" % (Alu,blocks,blocks,ParaStyle)
+    else:
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopLocal(queue,(jobs,),(jobs,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/jobs)),
+                                          numpy.uint32(nprnd(2**32/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)
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsNew)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+
+  circleCL.release()
+
+  #print jobs,numpy.mean(MyPi),numpy.median(MyPi),numpy.std(MyPi)
+  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):
+
+  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[2],
+                                                            coeffs_Mylq[3])
+    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^2+%.6f/N)" % \
+        (coeffs_Mylq2[0],coeffs_Mylq2[1],coeffs_Mylq2[2],coeffs_Mylq2[3],
+         coeffs_Mylq2[4])
+
+  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 or GPU
+  Alu='CPU'
+  # Id of GPU
+  Device=1
+  # GPU style can be Cuda (Nvidia implementation) or OpenCL
+  GpuStyle='OpenCL'
+  # Parallel distribution can be on Threads or Blocks
+  ParaStyle='Threads'
+  # Iterations is integer
+  Iterations=1000000
+  # JobStart in first number of Jobs to explore
+  JobStart=1
+  # JobEnd is last number of Jobs to explore
+  JobEnd=512
+  # 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
+  # Curves is True to print the curves
+  Curves=False
+
+  try:
+    opts, args = getopt.getopt(sys.argv[1:],"hoca:g:p:i:s:e:r:d:",["alu=","gpustyle=","parastyle=","iterations=","jobstart=","jobend=","redo=","device="])
+  except getopt.GetoptError:
+    print '%s -o (Out of Core Metrology) -c (Print Curves) -a <CPU/GPU> -d <DeviceId> -g <CUDA/OpenCL> -p <Threads/Hybrid/Blocks> -i <Iterations> -s <JobStart> -e <JobEnd> -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) -a <CPU/GPU> -d <DeviceId> -g <CUDA/OpenCL> -p <Threads/Hybrid/Blocks> -i <Iterations> -s <JobStart> -e <JobEnd> -r <RedoToImproveStats>' % sys.argv[0]
+      sys.exit()
+    elif opt == '-o':
+      OutMetrology=True
+    elif opt == '-c':
+      Curves=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.uint32(arg)
+    elif opt in ("-s", "--jobstart"):
+      JobStart = int(arg)
+    elif opt in ("-e", "--jobend"):
+      JobEnd = int(arg)
+    elif opt in ("-r", "--redo"):
+      Redo = int(arg)
+
+  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='Threads'
+
+  print "Compute unit : %s" % Alu
+  print "Device Identification : %s" % Device
+  print "GpuStyle used : %s" % GpuStyle
+  print "Parallel Style used : %s" % ParaStyle
+  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':
+    # For PyCUDA import
+    import pycuda.driver as cuda
+    import pycuda.gpuarray as gpuarray
+    import pycuda.autoinit
+    from pycuda.compiler import SourceModule
+
+  if GpuStyle=='OpenCL':
+    # 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)
+        Id=Id+1
+
+  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.uint32)
+
+    if OutMetrology:
+      duration=numpy.array([]).astype(numpy.float32)
+      for i in range(Redo):
+        start=time.time()
+        if GpuStyle=='CUDA':
+          try:
+            MetropolisCuda(circle,Iterations,1,Jobs,ParaStyle)
+          except:
+            print "Problem with %i // computations on Cuda" % Jobs
+        elif GpuStyle=='OpenCL':
+          try:
+            MetropolisOpenCL(circle,Iterations,1,Jobs,ParaStyle,Alu,Device)
+          except:
+            print "Problem with %i // computations on OpenCL" % Jobs            
+        duration=numpy.append(duration,time.time()-start)
+      avg=numpy.mean(duration)
+      med=numpy.median(duration)
+      std=numpy.std(duration)
+    else:
+      if GpuStyle=='CUDA':
+        try:
+          avg,med,std=MetropolisCuda(circle,Iterations,Redo,Jobs,ParaStyle)
+        except:
+          print "Problem with %i // computations on Cuda" % Jobs
+      elif GpuStyle=='OpenCL':
+        try:
+          avg,med,std=MetropolisOpenCL(circle,Iterations,Redo,Jobs,ParaStyle,Alu,Device)
+        except:
+          print "Problem with %i // computations on OpenCL" % Jobs            
+
+    if (avg,med,std) != (0,0,0):
+      print "avg,med,std",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
+    numpy.savez("Pi_%s_%s_%s_%s_%i_%.8i_%s" % (Alu,GpuStyle,ParaStyle,JobStart,JobEnd,Iterations,gethostname()),(ExploredJobs,average,median,stddev))
+    Jobs+=1
+
+  FitAndPrint(ExploredJobs,median,Curves)
+
Index: Pi/GPU/Pi-GPU.py-20130213
===================================================================
--- Pi/GPU/Pi-GPU.py-20130213	(revision 0)
+++ Pi/GPU/Pi-GPU.py-20130213	(revision 7)
@@ -0,0 +1,595 @@
+#!/usr/bin/env python
+
+#
+# Pi-by-MC using PyCUDA
+#
+# 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
+# nvidia-smi -c 1 : ko
+# nvidia-smi -c 3 : ko
+# nvidia-smi --gom=1 : not supported
+# 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 matplotlib.pyplot as plt
+import math
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+  return (T1*(s+c1*N+c2*N*N+p/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 + 2147483648.0f) * 2.328306435454494e-10f
+//#define KISSfp (KISS + 2147483648.0f) * 2.328306435454494e-10f
+#define MWCfp MWC * 2.328306435454494e-10f
+#define KISSfp KISS * 2.328306435454494e-10f
+
+__global__ void MainLoopBlocks(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(blockIdx.x+1);
+   uint w=seed_w/(blockIdx.x+1);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopThreads(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(threadIdx.x+1);
+   uint w=seed_w/(threadIdx.x+1);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[threadIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopHybrid(uint *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockDim.x*blockIdx.x+threadIdx.x]=total;
+   __syncthreads();
+
+}
+"""
+
+KERNEL_CODE_OPENCL="""
+
+// 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 + 2147483648.0f) * 2.328306435454494e-10f
+//#define KISSfp (KISS + 2147483648.0f) * 2.328306435454494e-10f
+#define MWCfp MWC * 2.328306435454494e-10f
+#define KISSfp KISS * 2.328306435454494e-10f
+
+__kernel void MainLoopGlobal(__global uint *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+   
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+     float y=MWCfp ;
+
+      // Matching test
+      int 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,threads,ParaStyle):
+
+  # Avec PyCUDA autoinit, rien a faire !
+  
+  circleCU = cuda.InOut(circle)
+  
+  mod = SourceModule(KERNEL_CODE_CUDA)
+
+  MetropolisBlocksCU=mod.get_function("MainLoopBlocks")
+  MetropolisThreadsCU=mod.get_function("MainLoopThreads")
+  MetropolisHybridCU=mod.get_function("MainLoopHybrid")
+  
+  start = pycuda.driver.Event()
+  stop = pycuda.driver.Event()
+  
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%threads==0:
+    iterationsCL=numpy.uint32(iterations/threads+1)
+    iterationsNew=iterationsCL*threads
+  else:
+    iterationsCL=numpy.uint32(iterations/threads)
+    iterationsNew=iterations
+
+  for i in range(steps):
+    start.record()
+    start.synchronize()
+    if ParaStyle=='Blocks':
+      MetropolisBlocksCU(circleCU,
+                         numpy.uint32(iterationsCL),
+                         numpy.uint32(nprnd(2**32/threads)),
+                         numpy.uint32(nprnd(2**32/threads)),
+                         grid=(threads,1),
+                         block=(1,1,1))
+      print "GPU with %i %s done" % (threads,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      blocks=threads/int(math.sqrt(float(threads)))
+      MetropolisHybridCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          grid=(blocks,1),
+                          block=(threads/blocks,1,1))
+      print "GPU with (blocks,threads)=(%i,%i) %s done" % (blocks,threads/blocks,ParaStyle)
+    else:
+      MetropolisThreadsCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          grid=(1,1),
+                          block=(threads,1,1))
+      print "GPU with %i %s done" % (threads,ParaStyle)
+    stop.record()
+    stop.synchronize()
+                
+    #elapsed = stop.time_since(start)*1e-3
+    elapsed = start.time_till(stop)*1e-3
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsCL)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+    #time.sleep(1)
+
+  print threads,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)
+
+  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+
+def MetropolisOpenCL(circle,iterations,steps,threads,ParaStyle,Alu):
+	
+  # Initialisation des variables en les CASTant correctement
+    
+  # Je detecte un peripherique GPU dans la liste des peripheriques
+  # for platform in cl.get_platforms():
+  # 	for device in platform.get_devices():
+  # 		if cl.device_type.to_string(device.type)=='GPU':
+  # 			GPU=device
+  #print "GPU detected: ",device.name
+
+  HasGPU=False
+  # Device selection based on choice (default is GPU)
+  for platform in cl.get_platforms():
+    for device in platform.get_devices():
+      if not HasGPU:
+        deviceType=cl.device_type.to_string(device.type)
+        if deviceType=="GPU" and Alu=="GPU":
+          GPU=device
+          print "GPU selected: ",device.name
+          HasGPU=True
+        if deviceType=="CPU" and Alu=="CPU":	    
+          GPU=device
+          print "CPU selected: ",device.name
+          HasGPU=True	
+				
+  # Je cree le contexte et la queue pour son execution
+  #ctx = cl.create_some_context()
+  ctx = cl.Context([GPU])
+  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")
+
+  #MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build()
+
+  i=0
+
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%threads==0:
+    iterationsCL=numpy.uint32(iterations/threads+1)
+    iterationsNew=iterationsCL*threads
+  else:
+    iterationsCL=numpy.uint32(iterations/threads)
+    iterationsNew=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
+      CLLaunch=MetropolisCL.MainLoopGlobal(queue,(threads,),None,
+                                           circleCL,
+                                           numpy.uint32(iterationsCL),
+                                           numpy.uint32(nprnd(2**32/threads)),
+                                           numpy.uint32(nprnd(2**32/threads)))
+      print "%s with %i %s done" % (Alu,threads,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      blocks=int(math.sqrt(threads))
+      CLLaunch=MetropolisCL.MainLoopHybrid(queue,(threads,),(blocks,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/threads)),
+                                          numpy.uint32(nprnd(2**32/threads)))
+      print "%s with (blocks,threads)=(%i,%i) %s done" % (Alu,blocks,threads/blocks,ParaStyle)
+    else:
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopLocal(queue,(threads,),(threads,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/threads)),
+                                          numpy.uint32(nprnd(2**32/threads)))
+      print "%s with %i %s done" % (Alu,threads,ParaStyle)
+
+    CLLaunch.wait()
+    cl.enqueue_copy(queue, circle, circleCL).wait()
+
+    elapsed = 1e-9*(CLLaunch.profile.end - CLLaunch.profile.start)
+
+    print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsNew)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+
+  circleCL.release()
+
+  #print threads,numpy.mean(MyPi),numpy.median(MyPi),numpy.std(MyPi)
+  print threads,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)
+	
+  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+
+def FitAndPrint(N,D,Curves):
+
+  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[2],
+                                                            coeffs_Mylq[3])
+    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^2+%.6f/N)" % \
+        (coeffs_Mylq2[0],coeffs_Mylq2[1],coeffs_Mylq2[2],coeffs_Mylq2[3],
+         coeffs_Mylq2[4])
+
+  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 or GPU
+  Alu='CPU'
+  # GPU style can be Cuda (Nvidia implementation) or OpenCL
+  GpuStyle='OpenCL'
+  # Parallel distribution can be on Threads or Blocks
+  ParaStyle='Threads'
+  # Iterations is integer
+  Iterations=1000000000
+  # ThreadStart in first number of Threads to explore
+  ThreadStart=1
+  # ThreadEnd is last number of Threads to explore
+  ThreadEnd=512
+  # 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
+  # Curves is True to print the curves
+  Curves=False
+
+  try:
+    opts, args = getopt.getopt(sys.argv[1:],"hoca:g:p:i:s:e:r:",["alu=","gpustyle=","parastyle=","iterations=","threadstart=","threadend=","redo="])
+  except getopt.GetoptError:
+    print '%s -o -a <CPU/GPU> -g <CUDA/OpenCL> -p <ParaStyle> -i <Iterations> -s <ThreadStart> -e <ThreadEnd> -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) -a <CPU/GPU> -g <CUDA/OpenCL> -p <Threads/Blocks> -i <Iterations> -s <ThreadStart> -e <ThreadEnd> -r <RedoToImproveStats>' % sys.argv[0]
+      sys.exit()
+    elif opt == '-o':
+      OutMetrology=True
+    elif opt == '-c':
+      Curves=True
+    elif opt in ("-a", "--alu"):
+      Alu = arg
+    elif opt in ("-g", "--gpustyle"):
+      GpuStyle = arg
+    elif opt in ("-p", "--parastyle"):
+      ParaStyle = arg
+    elif opt in ("-i", "--iterations"):
+      Iterations = numpy.uint32(arg)
+    elif opt in ("-s", "--threadstart"):
+      ThreadStart = int(arg)
+    elif opt in ("-e", "--threadend"):
+      ThreadEnd = int(arg)
+    elif opt in ("-r", "--redo"):
+      Redo = int(arg)
+
+  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='Threads'
+
+  print "Compute unit : %s" % Alu
+  print "GpuStyle used : %s" % GpuStyle
+  print "Parallel Style used : %s" % ParaStyle
+  print "Iterations : %s" % Iterations
+  print "Number of threads on start : %s" % ThreadStart
+  print "Number of threads on end : %s" % ThreadEnd
+  print "Number of redo : %s" % Redo
+  print "Metrology done out of CPU/GPU : %r" % OutMetrology
+
+  if GpuStyle=='CUDA':
+    # For PyCUDA import
+    import pycuda.driver as cuda
+    import pycuda.gpuarray as gpuarray
+    import pycuda.autoinit
+    from pycuda.compiler import SourceModule
+
+  if GpuStyle=='OpenCL':
+    # For PyOpenCL import
+    import pyopencl as cl
+
+  average=numpy.array([]).astype(numpy.float32)
+  median=numpy.array([]).astype(numpy.float32)
+  stddev=numpy.array([]).astype(numpy.float32)
+
+  ExploredThreads=numpy.array([]).astype(numpy.uint32)
+
+  Threads=ThreadStart
+
+  while Threads <= ThreadEnd:
+    ExploredThreads=numpy.append(ExploredThreads,Threads)
+    circle=numpy.zeros(Threads).astype(numpy.uint32)
+
+    if OutMetrology:
+      duration=numpy.array([]).astype(numpy.float32)
+      for i in range(Redo):
+        start=time.time()
+        if GpuStyle=='CUDA':
+          try:
+            MetropolisCuda(circle,Iterations,1,Threads,ParaStyle)
+          except:
+            print "Problem with %i // computations on Cuda" % Threads
+        elif GpuStyle=='OpenCL':
+          try:
+            MetropolisOpenCL(circle,Iterations,1,Threads,ParaStyle,Alu)
+          except:
+            print "Problem with %i // computations on OpenCL" % Threads            
+        duration=numpy.append(duration,time.time()-start)
+      avg=numpy.mean(duration)
+      med=numpy.median(duration)
+      std=numpy.std(duration)
+    else:
+      if GpuStyle=='CUDA':
+        avg,med,std=MetropolisCuda(circle,Iterations,Redo,Threads,ParaStyle)
+      elif GpuStyle=='OpenCL':
+        avg,med,std=MetropolisOpenCL(circle,Iterations,Redo,Threads,
+                                     ParaStyle,Alu)
+    print "avg,med,std",avg,med,std
+    average=numpy.append(average,avg)
+    median=numpy.append(median,med)
+    stddev=numpy.append(stddev,std)
+    #THREADS*=2
+    numpy.savez("Pi_%s_%s_%s_%s_%i_%.8i_%s" % (Alu,GpuStyle,ParaStyle,ThreadStart,ThreadEnd,Iterations,gethostname()),(ExploredThreads,average,median,stddev))
+    Threads+=1
+
+  FitAndPrint(ExploredThreads,median,Curves)
+
Index: Pi/GPU/Pi-GPU.py-20130205
===================================================================
--- Pi/GPU/Pi-GPU.py-20130205	(revision 0)
+++ Pi/GPU/Pi-GPU.py-20130205	(revision 7)
@@ -0,0 +1,519 @@
+#!/usr/bin/env python
+#
+# Pi-by-MC using PyCUDA
+#
+# 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
+# nvidia-smi -c 1 : ko
+# nvidia-smi -c 3 : ko
+# nvidia-smi --gom=1 : not supported
+# 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 matplotlib.pyplot as plt
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+  return (T1*(s+c1*N+c2*N*N+p/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 + 2147483648.0f) * 2.328306435454494e-10f
+//#define KISSfp (KISS + 2147483648.0f) * 2.328306435454494e-10f
+#define MWCfp MWC * 2.328306435454494e-10f
+#define KISSfp KISS * 2.328306435454494e-10f
+
+__global__ void MainLoopBlocks(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(blockIdx.x+1);
+   uint w=seed_w/(blockIdx.x+1);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopThreads(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(threadIdx.x+1);
+   uint w=seed_w/(threadIdx.x+1);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[threadIdx.x]=total;
+   __syncthreads();
+
+}
+"""
+
+KERNEL_CODE_OPENCL="""
+
+// 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 + 2147483648.0f) * 2.328306435454494e-10f
+//#define KISSfp (KISS + 2147483648.0f) * 2.328306435454494e-10f
+#define MWCfp MWC * 2.328306435454494e-10f
+#define KISSfp KISS * 2.328306435454494e-10f
+
+__kernel void MainLoopGlobal(__global int *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 int *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+   s[get_local_id(0)]=total;
+   barrier(CLK_LOCAL_MEM_FENCE);
+      
+}
+"""
+
+def MetropolisCuda(circle,iterations,steps,threads,ParaStyle):
+
+  # Avec PyCUDA autoinit, rien a faire !
+  
+  circleCU = cuda.InOut(circle)
+  
+  mod = SourceModule(KERNEL_CODE_CUDA)
+
+  MetropolisBlocksCU=mod.get_function("MainLoopBlocks")
+  MetropolisThreadsCU=mod.get_function("MainLoopThreads")
+  
+  start = pycuda.driver.Event()
+  stop = pycuda.driver.Event()
+  
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%threads==0:
+    iterationsCL=numpy.uint32(iterations/threads+1)
+    iterationsNew=iterationsCL*threads
+  else:
+    iterationsCL=numpy.uint32(iterations/threads)
+    iterationsNew=iterations
+
+  for i in range(steps):
+    start.record()
+    start.synchronize()
+    if ParaStyle=='Blocks':
+      MetropolisBlocksCU(circleCU,
+                         numpy.uint32(iterationsCL),
+                         numpy.uint32(nprnd(2**32/threads)),
+                         numpy.uint32(nprnd(2**32/threads)),
+                         grid=(threads,1),
+                         block=(1,1,1))
+    else:
+      MetropolisThreadsCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          grid=(1,1),
+                          block=(threads,1,1))
+    print "GPU done"
+    stop.record()
+    stop.synchronize()
+                
+    #elapsed = stop.time_since(start)*1e-3
+    elapsed = start.time_till(stop)*1e-3
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsCL)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+    #time.sleep(1)
+
+    print threads,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)
+
+  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+
+def MetropolisOpenCL(circle,iterations,steps,threads,ParaStyle,Alu):
+	
+  # Initialisation des variables en les CASTant correctement
+    
+  # Je detecte un peripherique GPU dans la liste des peripheriques
+  # for platform in cl.get_platforms():
+  # 	for device in platform.get_devices():
+  # 		if cl.device_type.to_string(device.type)=='GPU':
+  # 			GPU=device
+  #print "GPU detected: ",device.name
+
+  HasGPU=False
+  # Device selection based on choice (default is GPU)
+  for platform in cl.get_platforms():
+    for device in platform.get_devices():
+      if not HasGPU:
+        deviceType=cl.device_type.to_string(device.type)
+        if deviceType=="GPU" and Alu=="GPU":
+          GPU=device
+          print "GPU selected: ",device.name
+          HasGPU=True
+        if deviceType=="CPU" and Alu=="CPU":	    
+          GPU=device
+          print "CPU selected: ",device.name
+          HasGPU=True	
+				
+  # Je cree le contexte et la queue pour son execution
+  #ctx = cl.create_some_context()
+  ctx = cl.Context([GPU])
+  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")
+
+  #MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build()
+
+  i=0
+
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%threads==0:
+    iterationsCL=numpy.uint32(iterations/threads+1)
+    iterationsNew=iterationsCL*threads
+  else:
+    iterationsCL=numpy.uint32(iterations/threads)
+    iterationsNew=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
+      CLLaunch=MetropolisCL.MainLoopGlobal(queue,(threads,),(1,),
+                                           circleCL,
+                                           numpy.uint32(iterationsCL),
+                                           numpy.uint32(nprnd(2**32/threads)),
+                                           numpy.uint32(nprnd(2**32/threads)))
+      print "%s with %s done" % (Alu,ParaStyle)
+    else:
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopLocal(queue,(threads,),(threads,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/threads)),
+                                          numpy.uint32(nprnd(2**32/threads)))
+      print "%s with %s done" % (Alu,ParaStyle)
+
+    CLLaunch.wait()
+    cl.enqueue_copy(queue, circle, circleCL).wait()
+
+    elapsed = 1e-9*(CLLaunch.profile.end - CLLaunch.profile.start)
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsNew)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+
+  circleCL.release()
+
+  #print threads,numpy.mean(MyPi),numpy.median(MyPi),numpy.std(MyPi)
+  print threads,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)
+	
+  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+
+def FitAndPrint(N,D,Curves):
+
+  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(%.5f+%.5f/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(%.5f+%.5f/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(%.5f+%.5f*N+%.5f/N)" % (coeffs_Mylq[0],
+                                                            coeffs_Mylq[1],
+                                                            coeffs_Mylq[2],
+                                                            coeffs_Mylq[3])
+    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(%.5f+%.5f*N+%.5f*N^2+%.5f/N)" % \
+        (coeffs_Mylq2[0],coeffs_Mylq2[1],coeffs_Mylq2[2],coeffs_Mylq2[3],
+         coeffs_Mylq2[4])
+
+  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 or GPU
+  Alu='CPU'
+  # GPU style can be Cuda (Nvidia implementation) or OpenCL
+  GpuStyle='OpenCL'
+  # Parallel distribution can be on Threads or Blocks
+  ParaStyle='Threads'
+  # Iterations is integer
+  Iterations=1000000000
+  # ThreadStart in first number of Threads to explore
+  ThreadStart=1
+  # ThreadEnd is last number of Threads to explore
+  ThreadEnd=512
+  # 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
+  # Curves is True to print the curves
+  Curves=False
+
+  try:
+    opts, args = getopt.getopt(sys.argv[1:],"hoca:g:p:i:s:e:r:",["alu=","gpustyle=","parastyle=","iterations=","threadstart=","threadend=","redo="])
+  except getopt.GetoptError:
+    print '%s -o -a <CPU/GPU> -g <CUDA/OpenCL> -p <ParaStyle> -i <Iterations> -s <ThreadStart> -e <ThreadEnd> -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) -a <CPU/GPU> -g <CUDA/OpenCL> -p <Threads/Blocks> -i <Iterations> -s <ThreadStart> -e <ThreadEnd> -r <RedoToImproveStats>' % sys.argv[0]
+      sys.exit()
+    elif opt == '-o':
+      OutMetrology=True
+    elif opt == '-c':
+      Curves=True
+    elif opt in ("-a", "--alu"):
+      Alu = arg
+    elif opt in ("-g", "--gpustyle"):
+      GpuStyle = arg
+    elif opt in ("-p", "--parastyle"):
+      ParaStyle = arg
+    elif opt in ("-i", "--iterations"):
+      Iterations = numpy.uint32(arg)
+    elif opt in ("-s", "--threadstart"):
+      ThreadStart = int(arg)
+    elif opt in ("-e", "--threadend"):
+      ThreadEnd = int(arg)
+    elif opt in ("-r", "--redo"):
+      Redo = int(arg)
+
+  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'):
+    print "%s not exists, ParaStyle set as Threads !" % ParaStyle
+    ParaStyle='Threads'
+
+  print "Compute unit : %s" % Alu
+  print "GpuStyle used : %s" % GpuStyle
+  print "Parallel Style used : %s" % ParaStyle
+  print "Iterations : %s" % Iterations
+  print "Number of threads on start : %s" % ThreadStart
+  print "Number of threads on end : %s" % ThreadEnd
+  print "Number of redo : %s" % Redo
+  print "Metrology done out of CPU/GPU : %r" % OutMetrology
+
+  if GpuStyle=='CUDA':
+    # For PyCUDA import
+    import pycuda.driver as cuda
+    import pycuda.gpuarray as gpuarray
+    import pycuda.autoinit
+    from pycuda.compiler import SourceModule
+
+  if GpuStyle=='OpenCL':
+    # For PyOpenCL import
+    import pyopencl as cl
+
+  average=numpy.array([]).astype(numpy.float32)
+  median=numpy.array([]).astype(numpy.float32)
+  stddev=numpy.array([]).astype(numpy.float32)
+
+  ExploredThreads=numpy.array([]).astype(numpy.uint32)
+
+  Threads=ThreadStart
+
+  while Threads <= ThreadEnd:
+    ExploredThreads=numpy.append(ExploredThreads,Threads)
+    circle=numpy.zeros(Threads).astype(numpy.uint32)
+
+    if OutMetrology:
+      duration=numpy.array([]).astype(numpy.float32)
+      for i in range(Redo):
+        start=time.time()
+        if GpuStyle=='CUDA':
+          MetropolisCuda(circle,Iterations,1,Threads,ParaStyle)
+        elif GpuStyle=='OpenCL':
+          MetropolisOpenCL(circle,Iterations,1,Threads,ParaStyle,Alu)
+        duration=numpy.append(duration,time.time()-start)
+      avg=numpy.mean(duration)
+      med=numpy.median(duration)
+      std=numpy.std(duration)
+    else:
+      if GpuStyle=='CUDA':
+        avg,med,std=MetropolisCuda(circle,Iterations,Redo,Threads,ParaStyle)
+      elif GpuStyle=='OpenCL':
+        avg,med,std=MetropolisOpenCL(circle,Iterations,Redo,Threads,
+                                     ParaStyle,Alu)
+    print "avg,med,std",avg,med,std
+    average=numpy.append(average,avg)
+    median=numpy.append(median,med)
+    stddev=numpy.append(stddev,std)
+    #THREADS*=2
+    Threads+=1
+
+  numpy.savez("Pi_%s_%s_%s_%s_%i_%.8i_%s" % (Alu,GpuStyle,ParaStyle,ThreadStart,ThreadEnd,Iterations,gethostname()),(ExploredThreads,average,median,stddev))
+
+  FitAndPrint(ExploredThreads,median,Curves)
+
Index: Pi/GPU/Pi-GPU.py-20130214
===================================================================
--- Pi/GPU/Pi-GPU.py-20130214	(revision 0)
+++ Pi/GPU/Pi-GPU.py-20130214	(revision 7)
@@ -0,0 +1,602 @@
+#!/usr/bin/env python
+
+#
+# Pi-by-MC 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 matplotlib.pyplot as plt
+import math
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+  return (T1*(s+c1*N+c2*N*N+p/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(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(blockIdx.x+1);
+   uint w=seed_w/(blockIdx.x+1);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopThreads(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(threadIdx.x+1);
+   uint w=seed_w/(threadIdx.x+1);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[threadIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopHybrid(uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockDim.x*blockIdx.x+threadIdx.x]=total;
+   __syncthreads();
+
+}
+"""
+
+KERNEL_CODE_OPENCL="""
+
+// 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 uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+   
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+     float y=MWCfp ;
+
+      // Matching test
+      int 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):
+
+  # 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")
+  
+  start = pycuda.driver.Event()
+  stop = pycuda.driver.Event()
+  
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%jobs==0:
+    iterationsCL=numpy.uint32(iterations/jobs+1)
+    iterationsNew=iterationsCL*jobs
+  else:
+    iterationsCL=numpy.uint32(iterations/jobs)
+    iterationsNew=iterations
+
+  for i in range(steps):
+    start.record()
+    start.synchronize()
+    if ParaStyle=='Blocks':
+      MetropolisBlocksCU(circleCU,
+                         numpy.uint32(iterationsCL),
+                         numpy.uint32(nprnd(2**32/jobs)),
+                         numpy.uint32(nprnd(2**32/jobs)),
+                         grid=(jobs,1),
+                         block=(1,1,1))
+      print "GPU with %i %s done" % (jobs,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      blocks=jobs/int(math.sqrt(float(jobs)))
+      MetropolisHybridCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          grid=(blocks,1),
+                          block=(jobs/blocks,1,1))
+      print "GPU with (blocks,jobs)=(%i,%i) %s done" % (blocks,jobs/blocks,ParaStyle)
+    else:
+      MetropolisJobsCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          grid=(1,1),
+                          block=(jobs,1,1))
+      print "GPU with %i %s done" % (jobs,ParaStyle)
+    stop.record()
+    stop.synchronize()
+                
+    #elapsed = stop.time_since(start)*1e-3
+    elapsed = start.time_till(stop)*1e-3
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsCL)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+    #time.sleep(1)
+
+  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):
+	
+  # Initialisation des variables en les CASTant correctement
+    
+  # Je detecte un peripherique GPU dans la liste des peripheriques
+  # for platform in cl.get_platforms():
+  # 	for device in platform.get_devices():
+  # 		if cl.device_type.to_string(device.type)=='GPU':
+  # 			GPU=device
+  #print "GPU detected: ",device.name
+
+  HasGPU=False
+  Id=1
+  # Device selection based on choice (default is GPU)
+  for platform in cl.get_platforms():
+    for device in platform.get_devices():
+      if not HasGPU:
+        deviceType=cl.device_type.to_string(device.type)
+        if deviceType=="GPU" and Alu=="GPU" and Id==Device:
+          GPU=device
+          print "GPU selected: ",device.name
+          HasGPU=True
+        if deviceType=="CPU" and Alu=="CPU":	    
+          GPU=device
+          print "CPU selected: ",device.name
+          HasGPU=True
+      Id=Id+1
+				
+  # Je cree le contexte et la queue pour son execution
+  #ctx = cl.create_some_context()
+  ctx = cl.Context([GPU])
+  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")
+
+  #MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build()
+
+  i=0
+
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%jobs==0:
+    iterationsCL=numpy.uint32(iterations/jobs+1)
+    iterationsNew=iterationsCL*jobs
+  else:
+    iterationsCL=numpy.uint32(iterations/jobs)
+    iterationsNew=iterations
+
+  blocks=int(math.sqrt(jobs))
+
+  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
+      CLLaunch=MetropolisCL.MainLoopGlobal(queue,(jobs,),None,
+                                           circleCL,
+                                           numpy.uint32(iterationsCL),
+                                           numpy.uint32(nprnd(2**32/jobs)),
+                                           numpy.uint32(nprnd(2**32/jobs)))
+      print "%s with %i %s done" % (Alu,jobs,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopHybrid(queue,(blocks*blocks,),(blocks,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/jobs)),
+                                          numpy.uint32(nprnd(2**32/jobs)))
+      print "%s with (Blocks,Threads)=(%i,%i) %s done" % (Alu,blocks,blocks,ParaStyle)
+    else:
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopLocal(queue,(jobs,),(jobs,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/jobs)),
+                                          numpy.uint32(nprnd(2**32/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)
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsNew)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+
+  circleCL.release()
+
+  #print jobs,numpy.mean(MyPi),numpy.median(MyPi),numpy.std(MyPi)
+  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):
+
+  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[2],
+                                                            coeffs_Mylq[3])
+    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^2+%.6f/N)" % \
+        (coeffs_Mylq2[0],coeffs_Mylq2[1],coeffs_Mylq2[2],coeffs_Mylq2[3],
+         coeffs_Mylq2[4])
+
+  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 or GPU
+  Alu='CPU'
+  # Id of GPU
+  Device=1
+  # GPU style can be Cuda (Nvidia implementation) or OpenCL
+  GpuStyle='OpenCL'
+  # Parallel distribution can be on Threads or Blocks
+  ParaStyle='Threads'
+  # Iterations is integer
+  Iterations=1000000
+  # JobStart in first number of Jobs to explore
+  JobStart=1
+  # JobEnd is last number of Jobs to explore
+  JobEnd=512
+  # 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
+  # Curves is True to print the curves
+  Curves=False
+
+  try:
+    opts, args = getopt.getopt(sys.argv[1:],"hoca:g:p:i:s:e:r:d:",["alu=","gpustyle=","parastyle=","iterations=","jobstart=","jobend=","redo=","device="])
+  except getopt.GetoptError:
+    print '%s -o (Out of Core Metrology) -c (Print Curves) -a <CPU/GPU> -d <DeviceId> -g <CUDA/OpenCL> -p <Threads/Hybrid/Blocks> -i <Iterations> -s <JobStart> -e <JobEnd> -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) -a <CPU/GPU> -d <DeviceId> -g <CUDA/OpenCL> -p <Threads/Hybrid/Blocks> -i <Iterations> -s <JobStart> -e <JobEnd> -r <RedoToImproveStats>' % sys.argv[0]
+      sys.exit()
+    elif opt == '-o':
+      OutMetrology=True
+    elif opt == '-c':
+      Curves=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.uint32(arg)
+    elif opt in ("-s", "--jobstart"):
+      JobStart = int(arg)
+    elif opt in ("-e", "--jobend"):
+      JobEnd = int(arg)
+    elif opt in ("-r", "--redo"):
+      Redo = int(arg)
+
+  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='Threads'
+
+  print "Compute unit : %s" % Alu
+  print "Device Identification : %s" % Device
+  print "GpuStyle used : %s" % GpuStyle
+  print "Parallel Style used : %s" % ParaStyle
+  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':
+    # For PyCUDA import
+    import pycuda.driver as cuda
+    import pycuda.gpuarray as gpuarray
+    import pycuda.autoinit
+    from pycuda.compiler import SourceModule
+
+  if GpuStyle=='OpenCL':
+    # 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)
+        Id=Id+1
+
+  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.uint32)
+
+    if OutMetrology:
+      duration=numpy.array([]).astype(numpy.float32)
+      for i in range(Redo):
+        start=time.time()
+        if GpuStyle=='CUDA':
+          try:
+            MetropolisCuda(circle,Iterations,1,Jobs,ParaStyle)
+          except:
+            print "Problem with %i // computations on Cuda" % Jobs
+        elif GpuStyle=='OpenCL':
+          try:
+            MetropolisOpenCL(circle,Iterations,1,Jobs,ParaStyle,Alu,Device)
+          except:
+            print "Problem with %i // computations on OpenCL" % Jobs            
+        duration=numpy.append(duration,time.time()-start)
+      avg=numpy.mean(duration)
+      med=numpy.median(duration)
+      std=numpy.std(duration)
+    else:
+      if GpuStyle=='CUDA':
+        try:
+          avg,med,std=MetropolisCuda(circle,Iterations,Redo,Jobs,ParaStyle)
+        except:
+          print "Problem with %i // computations on Cuda" % Jobs
+      elif GpuStyle=='OpenCL':
+        try:
+          avg,med,std=MetropolisOpenCL(circle,Iterations,Redo,Jobs,ParaStyle,Alu,Device)
+        except:
+          print "Problem with %i // computations on OpenCL" % Jobs            
+
+    if (avg,med,std) != (0,0,0):
+      print "avg,med,std",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
+    numpy.savez("Pi_%s_%s_%s_%s_%i_%.8i_%s" % (Alu,GpuStyle,ParaStyle,JobStart,JobEnd,Iterations,gethostname()),(ExploredJobs,average,median,stddev))
+    Jobs+=1
+
+  FitAndPrint(ExploredJobs,median,Curves)
+
Index: Pi/GPU/Pi-GPU.py-20130206
===================================================================
--- Pi/GPU/Pi-GPU.py-20130206	(revision 0)
+++ Pi/GPU/Pi-GPU.py-20130206	(revision 7)
@@ -0,0 +1,524 @@
+#!/usr/bin/env python
+#
+# Pi-by-MC using PyCUDA
+#
+# 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
+# nvidia-smi -c 1 : ko
+# nvidia-smi -c 3 : ko
+# nvidia-smi --gom=1 : not supported
+# 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 matplotlib.pyplot as plt
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+  return (T1*(s+c1*N+c2*N*N+p/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 + 2147483648.0f) * 2.328306435454494e-10f
+//#define KISSfp (KISS + 2147483648.0f) * 2.328306435454494e-10f
+#define MWCfp MWC * 2.328306435454494e-10f
+#define KISSfp KISS * 2.328306435454494e-10f
+
+__global__ void MainLoopBlocks(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(blockIdx.x+1);
+   uint w=seed_w/(blockIdx.x+1);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopThreads(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(threadIdx.x+1);
+   uint w=seed_w/(threadIdx.x+1);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[threadIdx.x]=total;
+   __syncthreads();
+
+}
+"""
+
+KERNEL_CODE_OPENCL="""
+
+// 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 + 2147483648.0f) * 2.328306435454494e-10f
+//#define KISSfp (KISS + 2147483648.0f) * 2.328306435454494e-10f
+#define MWCfp MWC * 2.328306435454494e-10f
+#define KISSfp KISS * 2.328306435454494e-10f
+
+__kernel void MainLoopGlobal(__global int *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 int *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+   s[get_local_id(0)]=total;
+   barrier(CLK_LOCAL_MEM_FENCE);
+      
+}
+"""
+
+def MetropolisCuda(circle,iterations,steps,threads,ParaStyle):
+
+  # Avec PyCUDA autoinit, rien a faire !
+  
+  circleCU = cuda.InOut(circle)
+  
+  mod = SourceModule(KERNEL_CODE_CUDA)
+
+  MetropolisBlocksCU=mod.get_function("MainLoopBlocks")
+  MetropolisThreadsCU=mod.get_function("MainLoopThreads")
+  
+  start = pycuda.driver.Event()
+  stop = pycuda.driver.Event()
+  
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%threads==0:
+    iterationsCL=numpy.uint32(iterations/threads+1)
+    iterationsNew=iterationsCL*threads
+  else:
+    iterationsCL=numpy.uint32(iterations/threads)
+    iterationsNew=iterations
+
+  for i in range(steps):
+    start.record()
+    start.synchronize()
+    if ParaStyle=='Blocks':
+      MetropolisBlocksCU(circleCU,
+                         numpy.uint32(iterationsCL),
+                         numpy.uint32(nprnd(2**32/threads)),
+                         numpy.uint32(nprnd(2**32/threads)),
+                         grid=(threads,1),
+                         block=(1,1,1))
+    else:
+      MetropolisThreadsCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          grid=(1,1),
+                          block=(threads,1,1))
+    print "GPU done"
+    stop.record()
+    stop.synchronize()
+                
+    #elapsed = stop.time_since(start)*1e-3
+    elapsed = start.time_till(stop)*1e-3
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsCL)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+    #time.sleep(1)
+
+  print threads,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)
+
+  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+
+def MetropolisOpenCL(circle,iterations,steps,threads,ParaStyle,Alu):
+	
+  # Initialisation des variables en les CASTant correctement
+    
+  # Je detecte un peripherique GPU dans la liste des peripheriques
+  # for platform in cl.get_platforms():
+  # 	for device in platform.get_devices():
+  # 		if cl.device_type.to_string(device.type)=='GPU':
+  # 			GPU=device
+  #print "GPU detected: ",device.name
+
+  HasGPU=False
+  # Device selection based on choice (default is GPU)
+  for platform in cl.get_platforms():
+    for device in platform.get_devices():
+      if not HasGPU:
+        deviceType=cl.device_type.to_string(device.type)
+        if deviceType=="GPU" and Alu=="GPU":
+          GPU=device
+          print "GPU selected: ",device.name
+          HasGPU=True
+        if deviceType=="CPU" and Alu=="CPU":	    
+          GPU=device
+          print "CPU selected: ",device.name
+          HasGPU=True	
+				
+  # Je cree le contexte et la queue pour son execution
+  #ctx = cl.create_some_context()
+  ctx = cl.Context([GPU])
+  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")
+
+  #MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build()
+
+  i=0
+
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%threads==0:
+    iterationsCL=numpy.uint32(iterations/threads+1)
+    iterationsNew=iterationsCL*threads
+  else:
+    iterationsCL=numpy.uint32(iterations/threads)
+    iterationsNew=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
+      CLLaunch=MetropolisCL.MainLoopGlobal(queue,(threads,),(1,),
+                                           circleCL,
+                                           numpy.uint32(iterationsCL),
+                                           numpy.uint32(nprnd(2**32/threads)),
+                                           numpy.uint32(nprnd(2**32/threads)))
+      print "%s with %s done" % (Alu,ParaStyle)
+    else:
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopLocal(queue,(threads,),(threads,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/threads)),
+                                          numpy.uint32(nprnd(2**32/threads)))
+      print "%s with %s done" % (Alu,ParaStyle)
+
+    CLLaunch.wait()
+    cl.enqueue_copy(queue, circle, circleCL).wait()
+
+    elapsed = 1e-9*(CLLaunch.profile.end - CLLaunch.profile.start)
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsNew)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+
+  circleCL.release()
+
+  #print threads,numpy.mean(MyPi),numpy.median(MyPi),numpy.std(MyPi)
+  print threads,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)
+	
+  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+
+def FitAndPrint(N,D,Curves):
+
+  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(%.5f+%.5f/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(%.5f+%.5f/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(%.5f+%.5f*N+%.5f/N)" % (coeffs_Mylq[0],
+                                                            coeffs_Mylq[1],
+                                                            coeffs_Mylq[2],
+                                                            coeffs_Mylq[3])
+    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(%.5f+%.5f*N+%.5f*N^2+%.5f/N)" % \
+        (coeffs_Mylq2[0],coeffs_Mylq2[1],coeffs_Mylq2[2],coeffs_Mylq2[3],
+         coeffs_Mylq2[4])
+
+  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 or GPU
+  Alu='CPU'
+  # GPU style can be Cuda (Nvidia implementation) or OpenCL
+  GpuStyle='OpenCL'
+  # Parallel distribution can be on Threads or Blocks
+  ParaStyle='Threads'
+  # Iterations is integer
+  Iterations=1000000000
+  # ThreadStart in first number of Threads to explore
+  ThreadStart=1
+  # ThreadEnd is last number of Threads to explore
+  ThreadEnd=512
+  # 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
+  # Curves is True to print the curves
+  Curves=False
+
+  try:
+    opts, args = getopt.getopt(sys.argv[1:],"hoca:g:p:i:s:e:r:",["alu=","gpustyle=","parastyle=","iterations=","threadstart=","threadend=","redo="])
+  except getopt.GetoptError:
+    print '%s -o -a <CPU/GPU> -g <CUDA/OpenCL> -p <ParaStyle> -i <Iterations> -s <ThreadStart> -e <ThreadEnd> -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) -a <CPU/GPU> -g <CUDA/OpenCL> -p <Threads/Blocks> -i <Iterations> -s <ThreadStart> -e <ThreadEnd> -r <RedoToImproveStats>' % sys.argv[0]
+      sys.exit()
+    elif opt == '-o':
+      OutMetrology=True
+    elif opt == '-c':
+      Curves=True
+    elif opt in ("-a", "--alu"):
+      Alu = arg
+    elif opt in ("-g", "--gpustyle"):
+      GpuStyle = arg
+    elif opt in ("-p", "--parastyle"):
+      ParaStyle = arg
+    elif opt in ("-i", "--iterations"):
+      Iterations = numpy.uint32(arg)
+    elif opt in ("-s", "--threadstart"):
+      ThreadStart = int(arg)
+    elif opt in ("-e", "--threadend"):
+      ThreadEnd = int(arg)
+    elif opt in ("-r", "--redo"):
+      Redo = int(arg)
+
+  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'):
+    print "%s not exists, ParaStyle set as Threads !" % ParaStyle
+    ParaStyle='Threads'
+
+  print "Compute unit : %s" % Alu
+  print "GpuStyle used : %s" % GpuStyle
+  print "Parallel Style used : %s" % ParaStyle
+  print "Iterations : %s" % Iterations
+  print "Number of threads on start : %s" % ThreadStart
+  print "Number of threads on end : %s" % ThreadEnd
+  print "Number of redo : %s" % Redo
+  print "Metrology done out of CPU/GPU : %r" % OutMetrology
+
+  if GpuStyle=='CUDA':
+    # For PyCUDA import
+    import pycuda.driver as cuda
+    import pycuda.gpuarray as gpuarray
+    import pycuda.autoinit
+    from pycuda.compiler import SourceModule
+
+  if GpuStyle=='OpenCL':
+    # For PyOpenCL import
+    import pyopencl as cl
+
+  average=numpy.array([]).astype(numpy.float32)
+  median=numpy.array([]).astype(numpy.float32)
+  stddev=numpy.array([]).astype(numpy.float32)
+
+  ExploredThreads=numpy.array([]).astype(numpy.uint32)
+
+  Threads=ThreadStart
+
+  while Threads <= ThreadEnd:
+    ExploredThreads=numpy.append(ExploredThreads,Threads)
+    circle=numpy.zeros(Threads).astype(numpy.uint32)
+
+    if OutMetrology:
+      duration=numpy.array([]).astype(numpy.float32)
+      for i in range(Redo):
+        start=time.time()
+        if GpuStyle=='CUDA':
+          try:
+            MetropolisCuda(circle,Iterations,1,Threads,ParaStyle)
+          except:
+            print "Problem with %i // computations on Cuda" % Threads
+        elif GpuStyle=='OpenCL':
+          try:
+            MetropolisOpenCL(circle,Iterations,1,Threads,ParaStyle,Alu)
+          except:
+            print "Problem with %i // computations on OpenCL" % Threads            
+        duration=numpy.append(duration,time.time()-start)
+      avg=numpy.mean(duration)
+      med=numpy.median(duration)
+      std=numpy.std(duration)
+    else:
+      if GpuStyle=='CUDA':
+        avg,med,std=MetropolisCuda(circle,Iterations,Redo,Threads,ParaStyle)
+      elif GpuStyle=='OpenCL':
+        avg,med,std=MetropolisOpenCL(circle,Iterations,Redo,Threads,
+                                     ParaStyle,Alu)
+    print "avg,med,std",avg,med,std
+    average=numpy.append(average,avg)
+    median=numpy.append(median,med)
+    stddev=numpy.append(stddev,std)
+    #THREADS*=2
+    numpy.savez("Pi_%s_%s_%s_%s_%i_%.8i_%s" % (Alu,GpuStyle,ParaStyle,ThreadStart,ThreadEnd,Iterations,gethostname()),(ExploredThreads,average,median,stddev))
+    Threads+=1
+
+  FitAndPrint(ExploredThreads,median,Curves)
+
Index: Pi/GPU/Pi-GPU.py-20130207
===================================================================
--- Pi/GPU/Pi-GPU.py-20130207	(revision 0)
+++ Pi/GPU/Pi-GPU.py-20130207	(revision 7)
@@ -0,0 +1,595 @@
+#!/usr/bin/env python
+
+#
+# Pi-by-MC using PyCUDA
+#
+# 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
+# nvidia-smi -c 1 : ko
+# nvidia-smi -c 3 : ko
+# nvidia-smi --gom=1 : not supported
+# 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 matplotlib.pyplot as plt
+import math
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+  return (T1*(s+c1*N+c2*N*N+p/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 + 2147483648.0f) * 2.328306435454494e-10f
+//#define KISSfp (KISS + 2147483648.0f) * 2.328306435454494e-10f
+#define MWCfp MWC * 2.328306435454494e-10f
+#define KISSfp KISS * 2.328306435454494e-10f
+
+__global__ void MainLoopBlocks(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(blockIdx.x+1);
+   uint w=seed_w/(blockIdx.x+1);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopThreads(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(threadIdx.x+1);
+   uint w=seed_w/(threadIdx.x+1);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[threadIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopHybrid(uint *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockDim.x*blockIdx.x+threadIdx.x]=total;
+   __syncthreads();
+
+}
+"""
+
+KERNEL_CODE_OPENCL="""
+
+// 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 + 2147483648.0f) * 2.328306435454494e-10f
+//#define KISSfp (KISS + 2147483648.0f) * 2.328306435454494e-10f
+#define MWCfp MWC * 2.328306435454494e-10f
+#define KISSfp KISS * 2.328306435454494e-10f
+
+__kernel void MainLoopGlobal(__global uint *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+   
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+     float y=MWCfp ;
+
+      // Matching test
+      int 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,threads,ParaStyle):
+
+  # Avec PyCUDA autoinit, rien a faire !
+  
+  circleCU = cuda.InOut(circle)
+  
+  mod = SourceModule(KERNEL_CODE_CUDA)
+
+  MetropolisBlocksCU=mod.get_function("MainLoopBlocks")
+  MetropolisThreadsCU=mod.get_function("MainLoopThreads")
+  MetropolisHybridCU=mod.get_function("MainLoopHybrid")
+  
+  start = pycuda.driver.Event()
+  stop = pycuda.driver.Event()
+  
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%threads==0:
+    iterationsCL=numpy.uint32(iterations/threads+1)
+    iterationsNew=iterationsCL*threads
+  else:
+    iterationsCL=numpy.uint32(iterations/threads)
+    iterationsNew=iterations
+
+  for i in range(steps):
+    start.record()
+    start.synchronize()
+    if ParaStyle=='Blocks':
+      MetropolisBlocksCU(circleCU,
+                         numpy.uint32(iterationsCL),
+                         numpy.uint32(nprnd(2**32/threads)),
+                         numpy.uint32(nprnd(2**32/threads)),
+                         grid=(threads,1),
+                         block=(1,1,1))
+      print "GPU with %i %s done" % (threads,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      blocks=int(math.sqrt(float(threads)))
+      MetropolisHybridCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          grid=(blocks,1),
+                          block=(threads/blocks,1,1))
+      print "GPU with (blocks,threads)=(%i,%i) %s done" % (blocks,threads/blocks,ParaStyle)
+    else:
+      MetropolisThreadsCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          numpy.uint32(nprnd(2**32/threads)),
+                          grid=(1,1),
+                          block=(threads,1,1))
+      print "GPU with %i %s done" % (threads,ParaStyle)
+    stop.record()
+    stop.synchronize()
+                
+    #elapsed = stop.time_since(start)*1e-3
+    elapsed = start.time_till(stop)*1e-3
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsCL)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+    #time.sleep(1)
+
+  print threads,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)
+
+  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+
+def MetropolisOpenCL(circle,iterations,steps,threads,ParaStyle,Alu):
+	
+  # Initialisation des variables en les CASTant correctement
+    
+  # Je detecte un peripherique GPU dans la liste des peripheriques
+  # for platform in cl.get_platforms():
+  # 	for device in platform.get_devices():
+  # 		if cl.device_type.to_string(device.type)=='GPU':
+  # 			GPU=device
+  #print "GPU detected: ",device.name
+
+  HasGPU=False
+  # Device selection based on choice (default is GPU)
+  for platform in cl.get_platforms():
+    for device in platform.get_devices():
+      if not HasGPU:
+        deviceType=cl.device_type.to_string(device.type)
+        if deviceType=="GPU" and Alu=="GPU":
+          GPU=device
+          print "GPU selected: ",device.name
+          HasGPU=True
+        if deviceType=="CPU" and Alu=="CPU":	    
+          GPU=device
+          print "CPU selected: ",device.name
+          HasGPU=True	
+				
+  # Je cree le contexte et la queue pour son execution
+  #ctx = cl.create_some_context()
+  ctx = cl.Context([GPU])
+  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")
+
+  #MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build()
+
+  i=0
+
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%threads==0:
+    iterationsCL=numpy.uint32(iterations/threads+1)
+    iterationsNew=iterationsCL*threads
+  else:
+    iterationsCL=numpy.uint32(iterations/threads)
+    iterationsNew=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
+      CLLaunch=MetropolisCL.MainLoopGlobal(queue,(threads,),None,
+                                           circleCL,
+                                           numpy.uint32(iterationsCL),
+                                           numpy.uint32(nprnd(2**32/threads)),
+                                           numpy.uint32(nprnd(2**32/threads)))
+      print "%s with %i %s done" % (Alu,threads,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      blocks=int(math.sqrt(threads))
+      CLLaunch=MetropolisCL.MainLoopHybrid(queue,(threads,),(blocks,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/threads)),
+                                          numpy.uint32(nprnd(2**32/threads)))
+      print "%s with (blocks,threads)=(%i,%i) %s done" % (Alu,blocks,threads/blocks,ParaStyle)
+    else:
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopLocal(queue,(threads,),(threads,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/threads)),
+                                          numpy.uint32(nprnd(2**32/threads)))
+      print "%s with %i %s done" % (Alu,threads,ParaStyle)
+
+    CLLaunch.wait()
+    cl.enqueue_copy(queue, circle, circleCL).wait()
+
+    elapsed = 1e-9*(CLLaunch.profile.end - CLLaunch.profile.start)
+
+    print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsNew)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+
+  circleCL.release()
+
+  #print threads,numpy.mean(MyPi),numpy.median(MyPi),numpy.std(MyPi)
+  print threads,numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration)
+	
+  return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+
+def FitAndPrint(N,D,Curves):
+
+  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(%.5f+%.5f/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(%.5f+%.5f/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(%.5f+%.5f*N+%.5f/N)" % (coeffs_Mylq[0],
+                                                            coeffs_Mylq[1],
+                                                            coeffs_Mylq[2],
+                                                            coeffs_Mylq[3])
+    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(%.5f+%.5f*N+%.5f*N^2+%.5f/N)" % \
+        (coeffs_Mylq2[0],coeffs_Mylq2[1],coeffs_Mylq2[2],coeffs_Mylq2[3],
+         coeffs_Mylq2[4])
+
+  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 or GPU
+  Alu='CPU'
+  # GPU style can be Cuda (Nvidia implementation) or OpenCL
+  GpuStyle='OpenCL'
+  # Parallel distribution can be on Threads or Blocks
+  ParaStyle='Threads'
+  # Iterations is integer
+  Iterations=1000000000
+  # ThreadStart in first number of Threads to explore
+  ThreadStart=1
+  # ThreadEnd is last number of Threads to explore
+  ThreadEnd=512
+  # 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
+  # Curves is True to print the curves
+  Curves=False
+
+  try:
+    opts, args = getopt.getopt(sys.argv[1:],"hoca:g:p:i:s:e:r:",["alu=","gpustyle=","parastyle=","iterations=","threadstart=","threadend=","redo="])
+  except getopt.GetoptError:
+    print '%s -o -a <CPU/GPU> -g <CUDA/OpenCL> -p <ParaStyle> -i <Iterations> -s <ThreadStart> -e <ThreadEnd> -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) -a <CPU/GPU> -g <CUDA/OpenCL> -p <Threads/Blocks> -i <Iterations> -s <ThreadStart> -e <ThreadEnd> -r <RedoToImproveStats>' % sys.argv[0]
+      sys.exit()
+    elif opt == '-o':
+      OutMetrology=True
+    elif opt == '-c':
+      Curves=True
+    elif opt in ("-a", "--alu"):
+      Alu = arg
+    elif opt in ("-g", "--gpustyle"):
+      GpuStyle = arg
+    elif opt in ("-p", "--parastyle"):
+      ParaStyle = arg
+    elif opt in ("-i", "--iterations"):
+      Iterations = numpy.uint32(arg)
+    elif opt in ("-s", "--threadstart"):
+      ThreadStart = int(arg)
+    elif opt in ("-e", "--threadend"):
+      ThreadEnd = int(arg)
+    elif opt in ("-r", "--redo"):
+      Redo = int(arg)
+
+  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='Threads'
+
+  print "Compute unit : %s" % Alu
+  print "GpuStyle used : %s" % GpuStyle
+  print "Parallel Style used : %s" % ParaStyle
+  print "Iterations : %s" % Iterations
+  print "Number of threads on start : %s" % ThreadStart
+  print "Number of threads on end : %s" % ThreadEnd
+  print "Number of redo : %s" % Redo
+  print "Metrology done out of CPU/GPU : %r" % OutMetrology
+
+  if GpuStyle=='CUDA':
+    # For PyCUDA import
+    import pycuda.driver as cuda
+    import pycuda.gpuarray as gpuarray
+    import pycuda.autoinit
+    from pycuda.compiler import SourceModule
+
+  if GpuStyle=='OpenCL':
+    # For PyOpenCL import
+    import pyopencl as cl
+
+  average=numpy.array([]).astype(numpy.float32)
+  median=numpy.array([]).astype(numpy.float32)
+  stddev=numpy.array([]).astype(numpy.float32)
+
+  ExploredThreads=numpy.array([]).astype(numpy.uint32)
+
+  Threads=ThreadStart
+
+  while Threads <= ThreadEnd:
+    ExploredThreads=numpy.append(ExploredThreads,Threads)
+    circle=numpy.zeros(Threads).astype(numpy.uint32)
+
+    if OutMetrology:
+      duration=numpy.array([]).astype(numpy.float32)
+      for i in range(Redo):
+        start=time.time()
+        if GpuStyle=='CUDA':
+          try:
+            MetropolisCuda(circle,Iterations,1,Threads,ParaStyle)
+          except:
+            print "Problem with %i // computations on Cuda" % Threads
+        elif GpuStyle=='OpenCL':
+          try:
+            MetropolisOpenCL(circle,Iterations,1,Threads,ParaStyle,Alu)
+          except:
+            print "Problem with %i // computations on OpenCL" % Threads            
+        duration=numpy.append(duration,time.time()-start)
+      avg=numpy.mean(duration)
+      med=numpy.median(duration)
+      std=numpy.std(duration)
+    else:
+      if GpuStyle=='CUDA':
+        avg,med,std=MetropolisCuda(circle,Iterations,Redo,Threads,ParaStyle)
+      elif GpuStyle=='OpenCL':
+        avg,med,std=MetropolisOpenCL(circle,Iterations,Redo,Threads,
+                                     ParaStyle,Alu)
+    print "avg,med,std",avg,med,std
+    average=numpy.append(average,avg)
+    median=numpy.append(median,med)
+    stddev=numpy.append(stddev,std)
+    #THREADS*=2
+    numpy.savez("Pi_%s_%s_%s_%s_%i_%.8i_%s" % (Alu,GpuStyle,ParaStyle,ThreadStart,ThreadEnd,Iterations,gethostname()),(ExploredThreads,average,median,stddev))
+    Threads+=1
+
+  FitAndPrint(ExploredThreads,median,Curves)
+
Index: Pi/GPU/Pi-GPU.py
===================================================================
--- Pi/GPU/Pi-GPU.py	(revision 0)
+++ Pi/GPU/Pi-GPU.py	(revision 7)
@@ -0,0 +1,602 @@
+#!/usr/bin/env python
+
+#
+# Pi-by-MC 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 matplotlib.pyplot as plt
+import math
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+  return (T1*(s+c1*N+c2*N*N+p/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(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(blockIdx.x+1);
+   uint w=seed_w/(blockIdx.x+1);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopThreads(uint *s,uint iterations,uint seed_w,uint seed_z)
+{
+   uint z=seed_z/(threadIdx.x+1);
+   uint w=seed_w/(threadIdx.x+1);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[threadIdx.x]=total;
+   __syncthreads();
+
+}
+
+__global__ void MainLoopHybrid(uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+
+   }
+
+   s[blockDim.x*blockIdx.x+threadIdx.x]=total;
+   __syncthreads();
+
+}
+"""
+
+KERNEL_CODE_OPENCL="""
+
+// 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 uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int 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 uint *s,uint 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);
+   
+   // uint jsr=123456789;
+   // uint jcong=380116160;
+
+   int total=0;
+
+   for (uint i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+     float y=MWCfp ;
+
+      // Matching test
+      int 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):
+
+  # 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")
+  
+  start = pycuda.driver.Event()
+  stop = pycuda.driver.Event()
+  
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%jobs==0:
+    iterationsCL=numpy.uint32(iterations/jobs)
+    iterationsNew=iterationsCL*jobs
+  else:
+    iterationsCL=numpy.uint32(iterations/jobs+1)
+    iterationsNew=iterations
+
+  for i in range(steps):
+    start.record()
+    start.synchronize()
+    if ParaStyle=='Blocks':
+      MetropolisBlocksCU(circleCU,
+                         numpy.uint32(iterationsCL),
+                         numpy.uint32(nprnd(2**32/jobs)),
+                         numpy.uint32(nprnd(2**32/jobs)),
+                         grid=(jobs,1),
+                         block=(1,1,1))
+      print "GPU with %i %s done" % (jobs,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      blocks=jobs/int(math.sqrt(float(jobs)))
+      MetropolisHybridCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          grid=(blocks,1),
+                          block=(jobs/blocks,1,1))
+      print "GPU with (blocks,jobs)=(%i,%i) %s done" % (blocks,jobs/blocks,ParaStyle)
+    else:
+      MetropolisJobsCU(circleCU,
+                          numpy.uint32(iterationsCL),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          numpy.uint32(nprnd(2**32/jobs)),
+                          grid=(1,1),
+                          block=(jobs,1,1))
+      print "GPU with %i %s done" % (jobs,ParaStyle)
+    stop.record()
+    stop.synchronize()
+                
+    #elapsed = stop.time_since(start)*1e-3
+    elapsed = start.time_till(stop)*1e-3
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsCL)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+    #time.sleep(1)
+
+  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):
+	
+  # Initialisation des variables en les CASTant correctement
+    
+  # Je detecte un peripherique GPU dans la liste des peripheriques
+  # for platform in cl.get_platforms():
+  # 	for device in platform.get_devices():
+  # 		if cl.device_type.to_string(device.type)=='GPU':
+  # 			GPU=device
+  #print "GPU detected: ",device.name
+
+  HasGPU=False
+  Id=1
+  # Device selection based on choice (default is GPU)
+  for platform in cl.get_platforms():
+    for device in platform.get_devices():
+      if not HasGPU:
+        deviceType=cl.device_type.to_string(device.type)
+        if deviceType=="GPU" and Alu=="GPU" and Id==Device:
+          GPU=device
+          print "GPU selected: ",device.name
+          HasGPU=True
+        if deviceType=="CPU" and Alu=="CPU":	    
+          GPU=device
+          print "CPU selected: ",device.name
+          HasGPU=True
+      Id=Id+1
+				
+  # Je cree le contexte et la queue pour son execution
+  #ctx = cl.create_some_context()
+  ctx = cl.Context([GPU])
+  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")
+
+  #MetropolisCL = cl.Program(ctx,KERNEL_CODE_OPENCL).build()
+
+  i=0
+
+  MyPi=numpy.zeros(steps)
+  MyDuration=numpy.zeros(steps)
+  
+  if iterations%jobs==0:
+    iterationsCL=numpy.uint32(iterations/jobs)
+    iterationsNew=iterationsCL*jobs
+  else:
+    iterationsCL=numpy.uint32(iterations/jobs+1)
+    iterationsNew=iterations
+
+  blocks=int(math.sqrt(jobs))
+
+  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
+      CLLaunch=MetropolisCL.MainLoopGlobal(queue,(jobs,),None,
+                                           circleCL,
+                                           numpy.uint32(iterationsCL),
+                                           numpy.uint32(nprnd(2**32/jobs)),
+                                           numpy.uint32(nprnd(2**32/jobs)))
+      print "%s with %i %s done" % (Alu,jobs,ParaStyle)
+    elif ParaStyle=='Hybrid':
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopHybrid(queue,(blocks*blocks,),(blocks,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/jobs)),
+                                          numpy.uint32(nprnd(2**32/jobs)))
+      print "%s with (Blocks,Threads)=(%i,%i) %s done" % (Alu,blocks,blocks,ParaStyle)
+    else:
+      # en OpenCL, necessaire de mettre un Global_id identique au local_id
+      CLLaunch=MetropolisCL.MainLoopLocal(queue,(jobs,),(jobs,),
+                                          circleCL,
+                                          numpy.uint32(iterationsCL),
+                                          numpy.uint32(nprnd(2**32/jobs)),
+                                          numpy.uint32(nprnd(2**32/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)
+
+    #print circle,float(numpy.sum(circle))
+    MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsNew)
+    MyDuration[i]=elapsed
+    #print MyPi[i],MyDuration[i]
+
+  circleCL.release()
+
+  #print jobs,numpy.mean(MyPi),numpy.median(MyPi),numpy.std(MyPi)
+  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):
+
+  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[2],
+                                                            coeffs_Mylq[3])
+    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^2+%.6f/N)" % \
+        (coeffs_Mylq2[0],coeffs_Mylq2[1],coeffs_Mylq2[2],coeffs_Mylq2[3],
+         coeffs_Mylq2[4])
+
+  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 or GPU
+  Alu='CPU'
+  # Id of GPU
+  Device=1
+  # 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
+  # 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
+  # Curves is True to print the curves
+  Curves=False
+
+  try:
+    opts, args = getopt.getopt(sys.argv[1:],"hoca:g:p:i:s:e:r:d:",["alu=","gpustyle=","parastyle=","iterations=","jobstart=","jobend=","redo=","device="])
+  except getopt.GetoptError:
+    print '%s -o (Out of Core Metrology) -c (Print Curves) -a <CPU/GPU> -d <DeviceId> -g <CUDA/OpenCL> -p <Threads/Hybrid/Blocks> -i <Iterations> -s <JobStart> -e <JobEnd> -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) -a <CPU/GPU> -d <DeviceId> -g <CUDA/OpenCL> -p <Threads/Hybrid/Blocks> -i <Iterations> -s <JobStart> -e <JobEnd> -r <RedoToImproveStats>' % sys.argv[0]
+      sys.exit()
+    elif opt == '-o':
+      OutMetrology=True
+    elif opt == '-c':
+      Curves=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.uint32(arg)
+    elif opt in ("-s", "--jobstart"):
+      JobStart = int(arg)
+    elif opt in ("-e", "--jobend"):
+      JobEnd = int(arg)
+    elif opt in ("-r", "--redo"):
+      Redo = int(arg)
+
+  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='Threads'
+
+  print "Compute unit : %s" % Alu
+  print "Device Identification : %s" % Device
+  print "GpuStyle used : %s" % GpuStyle
+  print "Parallel Style used : %s" % ParaStyle
+  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':
+    # For PyCUDA import
+    import pycuda.driver as cuda
+    import pycuda.gpuarray as gpuarray
+    import pycuda.autoinit
+    from pycuda.compiler import SourceModule
+
+  if GpuStyle=='OpenCL':
+    # 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)
+        Id=Id+1
+
+  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.uint32)
+
+    if OutMetrology:
+      duration=numpy.array([]).astype(numpy.float32)
+      for i in range(Redo):
+        start=time.time()
+        if GpuStyle=='CUDA':
+          try:
+            MetropolisCuda(circle,Iterations,1,Jobs,ParaStyle)
+          except:
+            print "Problem with %i // computations on Cuda" % Jobs
+        elif GpuStyle=='OpenCL':
+          try:
+            MetropolisOpenCL(circle,Iterations,1,Jobs,ParaStyle,Alu,Device)
+          except:
+            print "Problem with %i // computations on OpenCL" % Jobs            
+        duration=numpy.append(duration,time.time()-start)
+      avg=numpy.mean(duration)
+      med=numpy.median(duration)
+      std=numpy.std(duration)
+    else:
+      if GpuStyle=='CUDA':
+        try:
+          avg,med,std=MetropolisCuda(circle,Iterations,Redo,Jobs,ParaStyle)
+        except:
+          print "Problem with %i // computations on Cuda" % Jobs
+      elif GpuStyle=='OpenCL':
+        try:
+          avg,med,std=MetropolisOpenCL(circle,Iterations,Redo,Jobs,ParaStyle,Alu,Device)
+        except:
+          print "Problem with %i // computations on OpenCL" % Jobs            
+
+    if (avg,med,std) != (0,0,0):
+      print "avg,med,std",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
+    numpy.savez("Pi_%s_%s_%s_%s_%i_%.8i_%s" % (Alu,GpuStyle,ParaStyle,JobStart,JobEnd,Iterations,gethostname()),(ExploredJobs,average,median,stddev))
+    Jobs+=1
+
+  FitAndPrint(ExploredJobs,median,Curves)
+
Index: Pi/MP/PiMP-unprotected.py
===================================================================
--- Pi/MP/PiMP-unprotected.py	(revision 0)
+++ Pi/MP/PiMP-unprotected.py	(revision 7)
@@ -0,0 +1,222 @@
+#!/usr/bin/env python
+
+# Version using Multiprocessing module
+#
+# Pi-by-MC
+#
+# CC BY-NC-SA 2013 : <emmanuel.quemener@ens-lyon.fr> 
+#
+
+import getopt
+import numpy
+
+from random import random
+# Multithread library call
+from multiprocessing import Pool
+
+# Common tools
+import sys
+import getopt
+import time
+import matplotlib.pyplot as plt
+import math
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+    return (T1*(s+c1*N+c2*N*N+p/N))
+
+def MainLoop(iterations):
+    
+    total=0
+    # Absulute necessary to use xrange instead of range to avoid out of memory !
+    for i in xrange(iterations):
+        # Random access coordonate
+        x,y=random(),random()
+        
+        if ((x*x+y*y) < 1.0):
+            total+=1
+            
+    return(total)
+        
+def MetropolisMP(circle,iterations,steps,jobs):
+
+    MyPi=numpy.zeros(steps)
+    MyDuration=numpy.array([])
+    
+    # Define iterations to send to each node
+    if iterations%jobs==0:
+        iterationsMP=iterations/jobs
+    else:
+        iterationsMP=iterations/jobs+1
+    print "%i iterations will be send to each core" % iterationsMP
+    
+    WorkItems=[]
+    for i in xrange(0,jobs):
+        WorkItems.append(iterationsMP)
+
+    for i in xrange(steps):
+        start=time.time()
+        pool=Pool(processes=jobs)
+        # Define the number of processes to be launched at a time
+        # pool=Pool(processes=CORES)
+        print "Start on %i processors..." % jobs
+        # MAP: distribution of usecases T to be applied to MetropolisStrip 
+        # WorkLaunches=[ pool.apply_async(MainLoop,(wi,)) for wi in WorkItems ]
+        # circle=[wl.get() for wl in WorkLaunches]
+        circle=pool.map(MainLoop,WorkItems)
+        MyDuration=numpy.append(MyDuration,time.time()-start)
+        MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsMP)/float(jobs)
+
+    return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+def FitAndPrint(N,D,Curves):
+
+    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[2],
+                                                                coeffs_Mylq[3])
+        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) 
+
+    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
+    # Iterations is integer
+    Iterations=1000000
+    # JobStart in first number of Jobs to explore
+    JobStart=1
+    # JobEnd is last number of Jobs to explore
+    JobEnd=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
+    # Curves is True to print the curves
+    Curves=False
+
+    try:
+        opts, args = getopt.getopt(sys.argv[1:],"hoci:s:e:r:",["alu=","gpustyle=","parastyle=","iterations=","jobstart=","jobend=","redo=","device="])
+    except getopt.GetoptError:
+        print '%s -o (Out of Core Metrology) -c (Print Curves) -i <Iterations> -s <JobStart> -e <JobEnd> -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) -i <Iterations> -s <JobStart> -e <JobEnd> -r <RedoToImproveStats>' % sys.argv[0]
+            sys.exit()
+        elif opt == '-o':
+            OutMetrology=True
+        elif opt == '-c':
+            Curves=True
+        elif opt in ("-i", "--iterations"):
+            Iterations = numpy.uint32(arg)
+        elif opt in ("-s", "--jobstart"):
+            JobStart = int(arg)
+        elif opt in ("-e", "--jobend"):
+            JobEnd = int(arg)
+        elif opt in ("-r", "--redo"):
+            Redo = int(arg)
+
+    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 : %r" % OutMetrology
+
+    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.uint32)
+        
+        if OutMetrology:
+            duration=numpy.array([]).astype(numpy.float32)
+            for i in range(Redo):
+                start=time.time()
+                MetropolisMP(circle,Iterations,Redo,Jobs)
+                duration=numpy.append(duration,time.time()-start)
+            avg=numpy.mean(duration)
+            med=numpy.median(duration)
+            std=numpy.std(duration)
+        else:
+            avg,med,std=MetropolisMP(circle,Iterations,Redo,Jobs)
+
+        if (avg,med,std) != (0,0,0):
+            print "avg,med,std",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
+        numpy.savez("Pi_%s_%i_%.8i_%s" % (JobStart,JobEnd,Iterations,gethostname()),(ExploredJobs,average,median,stddev))
+        Jobs+=1
+
+    FitAndPrint(ExploredJobs,median,Curves)

Property changes on: Pi/MP/PiMP-unprotected.py
___________________________________________________________________
Added: svn:executable
   + *

Index: Pi/MP/PiMP.py.orig
===================================================================
--- Pi/MP/PiMP.py.orig	(revision 0)
+++ Pi/MP/PiMP.py.orig	(revision 7)
@@ -0,0 +1,61 @@
+#!/usr/bin/env python
+
+# If numpy is loaded without being used, it does not work !
+from math import exp
+from random import random
+# Multithread library call
+#import multiprocessing
+from multiprocessing import Pool
+import numpy
+from numpy.random import rand
+
+
+def MainLoop(iterations):
+
+    total=numpy.int(0)
+    for i in range(0,iterations):
+        # Random access coordonate
+        #x,y=random(),random()
+        x,y=rand(),rand()
+
+        if ((x*x+y*y) < 1.0):
+                total+=1
+
+    return(total)
+
+if __name__=='__main__':
+
+    # Machine
+    CORES=4
+    
+    # A=numpy.zeros(10)
+    
+    # Au dessus de 4 10^7, le MPI ne veut plus se lancer...
+    Iterations=100000000
+
+    total=0
+    
+    # Define iterations to send to each node
+    if Iterations%CORES==0:
+        iterations=Iterations/CORES
+    else:
+        iterations=Iterations/CORES+1
+        print "%i iterations will be send to each core" % iterations
+        
+    IT=[]
+    for i in range(0,CORES):
+        IT.append(iterations)
+        
+    print IT
+
+    # Define the number of processes to be launched at a time
+    # POOL: resources to be used ( 
+    # pool=Pool(processes=CORES)
+    pool=Pool(processes=CORES)
+    print "Start on %i processors..." % CORES
+
+    # MAP: distribution of usecases T to be applied to MetropolisStrip 
+    Results=pool.map(MainLoop,IT)
+    
+    # print sum(Results.get())
+    print Results

Property changes on: Pi/MP/PiMP.py.orig
___________________________________________________________________
Added: svn:executable
   + *

Index: Pi/MP/PostProcess.py
===================================================================
--- Pi/MP/PostProcess.py	(revision 0)
+++ Pi/MP/PostProcess.py	(revision 7)
@@ -0,0 +1,12 @@
+import numpy
+npzfile=numpy.load("Ising2D_MP_256_16777216.npz")
+npzfile['arr_0']
+T=npzfile['arr_0'].transpose()[0]
+E=npzfile['arr_0'].transpose()[1]
+M=npzfile['arr_0'].transpose()[2]
+print T
+print E
+print M
+import matplotlib.pyplot 
+matplotlib.pyplot.plot(T,E)
+matplotlib.pyplot.show()
Index: Pi/MP/PiMP.py
===================================================================
--- Pi/MP/PiMP.py	(revision 0)
+++ Pi/MP/PiMP.py	(revision 7)
@@ -0,0 +1,228 @@
+#!/usr/bin/env python
+
+# Version using Multiprocessing module
+#
+# Pi-by-MC
+#
+# CC BY-NC-SA 2013 : <emmanuel.quemener@ens-lyon.fr> 
+#
+
+import getopt
+import numpy
+
+from random import random
+# Multithread library call
+from multiprocessing import Pool
+
+# Common tools
+import sys
+import getopt
+import time
+import matplotlib.pyplot as plt
+import math
+from scipy.optimize import curve_fit
+from socket import gethostname
+
+# 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+c*N+p/N))
+
+# Predicted Mylq Law with second order
+def Mylq2(N, T1,s,c1,c2,p):
+    return (T1*(s+c1*N+c2*N*N+p/N))
+
+def MainLoop(iterations):
+    
+    total=0
+    # Absulute necessary to use xrange instead of range to avoid out of memory !
+    for i in xrange(iterations):
+        # Random access coordonate
+        x,y=random(),random()
+        
+        if ((x*x+y*y) < 1.0):
+            total+=1
+            
+    return(total)
+        
+def MetropolisMP(circle,iterations,steps,jobs):
+
+    MyPi=numpy.zeros(steps)
+    MyDuration=numpy.array([])
+    
+    # Define iterations to send to each node
+    if iterations%jobs==0:
+        iterationsMP=iterations/jobs
+    else:
+        iterationsMP=iterations/jobs+1
+    print "%i iterations will be send to each core" % iterationsMP
+    
+    WorkItems=[]
+    for i in xrange(0,jobs):
+        WorkItems.append(iterationsMP)
+
+    for i in xrange(steps):
+        start=time.time()
+        pool=Pool(processes=jobs)
+        # Define the number of processes to be launched at a time
+        # pool=Pool(processes=CORES)
+        print "Start on %i processors..." % jobs
+        # MAP: distribution of usecases T to be applied to MetropolisStrip 
+        # WorkLaunches=[ pool.apply_async(MainLoop,(wi,)) for wi in WorkItems ]
+        # circle=[wl.get() for wl in WorkLaunches]
+        circle=pool.map(MainLoop,WorkItems)
+        MyDuration=numpy.append(MyDuration,time.time()-start)
+        MyPi[i]=4.*float(numpy.sum(circle))/float(iterationsMP)/float(jobs)
+
+    return(numpy.mean(MyDuration),numpy.median(MyDuration),numpy.std(MyDuration))
+
+def FitAndPrint(N,D,Curves):
+
+    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[2],
+                                                                coeffs_Mylq[3])
+        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) 
+
+    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
+    # Iterations is integer
+    Iterations=1000000
+    # JobStart in first number of Jobs to explore
+    JobStart=1
+    # JobEnd is last number of Jobs to explore
+    JobEnd=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
+    # Curves is True to print the curves
+    Curves=False
+
+    try:
+        opts, args = getopt.getopt(sys.argv[1:],"hoci:s:e:r:",["alu=","gpustyle=","parastyle=","iterations=","jobstart=","jobend=","redo=","device="])
+    except getopt.GetoptError:
+        print '%s -o (Out of Core Metrology) -c (Print Curves) -i <Iterations> -s <JobStart> -e <JobEnd> -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) -i <Iterations> -s <JobStart> -e <JobEnd> -r <RedoToImproveStats>' % sys.argv[0]
+            sys.exit()
+        elif opt == '-o':
+            OutMetrology=True
+        elif opt == '-c':
+            Curves=True
+        elif opt in ("-i", "--iterations"):
+            Iterations = numpy.uint32(arg)
+        elif opt in ("-s", "--jobstart"):
+            JobStart = int(arg)
+        elif opt in ("-e", "--jobend"):
+            JobEnd = int(arg)
+        elif opt in ("-r", "--redo"):
+            Redo = int(arg)
+
+    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 : %r" % OutMetrology
+
+    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.uint32)
+        
+        if OutMetrology:
+            duration=numpy.array([]).astype(numpy.float32)
+            for i in range(Redo):
+                start=time.time()
+                try:
+                    MetropolisMP(circle,Iterations,Redo,Jobs)
+                except:
+                    print "Problem with %i // computations" % Jobs
+                duration=numpy.append(duration,time.time()-start)
+            avg=numpy.mean(duration)
+            med=numpy.median(duration)
+            std=numpy.std(duration)
+        else:
+            try:
+                avg,med,std=MetropolisMP(circle,Iterations,Redo,Jobs)
+            except:
+                print "Problem with %i // computations" % Jobs
+
+        if (avg,med,std) != (0,0,0):
+            print "avg,med,std",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
+        numpy.savez("Pi_%s_%i_%.8i_%s" % (JobStart,JobEnd,Iterations,gethostname()),(ExploredJobs,average,median,stddev))
+        Jobs+=1
+
+    FitAndPrint(ExploredJobs,median,Curves)

Property changes on: Pi/MP/PiMP.py
___________________________________________________________________
Added: svn:executable
   + *

Index: Pi/MP/PiMP.py.orig2
===================================================================
--- Pi/MP/PiMP.py.orig2	(revision 0)
+++ Pi/MP/PiMP.py.orig2	(revision 7)
@@ -0,0 +1,54 @@
+#!/usr/bin/env python
+
+# If numpy is loaded without being used, it does not work !
+from math import exp
+from random import random
+# Multithread library call
+from multiprocessing import Pool
+
+def MainLoop(iterations):
+
+    total=0
+    for i in xrange(iterations):
+        # Random access coordonate
+        x,y=random(),random()
+
+        if ((x*x+y*y) < 1.0):
+                total+=1
+
+    return(total)
+
+if __name__=='__main__':
+
+   # Machine
+   CORES=4
+
+   # Au dessus de 4 10^7, le MPI ne veut plus se lancer...
+   Iterations=1000000000
+
+   total=0
+
+   # Define iterations to send to each node
+   if Iterations%CORES==0:
+	   iterations=Iterations/CORES
+   else:
+	   iterations=Iterations/CORES+1
+   print "%i iterations will be send to each core" % iterations
+
+   IT=[]
+   for i in range(0,CORES):
+	   IT.append(iterations)
+
+   print IT
+
+   # Define the number of processes to be launched at a time
+   # POOL: resources to be used ( 
+   #pool=Pool(processes=CORES)
+   pool=Pool(CORES)
+   print "Start on %i processors..." % CORES
+
+   # MAP: distribution of usecases T to be applied to MetropolisStrip 
+   Results=[ pool.apply_async(MainLoop,(i,)) for i in IT]
+
+   results=[r.get() for r in Results]
+   print results,sum(results)

Property changes on: Pi/MP/PiMP.py.orig2
___________________________________________________________________
Added: svn:executable
   + *

Index: Pi/Serial/PostProcess.py
===================================================================
--- Pi/Serial/PostProcess.py	(revision 0)
+++ Pi/Serial/PostProcess.py	(revision 7)
@@ -0,0 +1,10 @@
+import numpy
+import sys
+npzfile=numpy.load(sys.argv[1])
+T=npzfile['arr_0'][0]
+E=npzfile['arr_0'][1]
+ED=npzfile['arr_0'][2]
+import matplotlib.pyplot
+matplotlib.pyplot.plot(T,E)
+#matplotlib.pyplot.plot(T,E,T,ED)
+matplotlib.pyplot.show()
Index: Pi/Results/PostProcessCGPU.py
===================================================================
--- Pi/Results/PostProcessCGPU.py	(revision 0)
+++ Pi/Results/PostProcessCGPU.py	(revision 7)
@@ -0,0 +1,49 @@
+import numpy
+import sys
+
+npzOCLLG=numpy.load("Ising2D_OCLL_256_16777216_GPU_atlas.npz")
+npzOCLLC=numpy.load("Ising2D_OCLL_256_16777216_CPU_atlas.npz")
+npzCUDAL=numpy.load("Ising2D_CUDAL_256_16777216_GPU_atlas.npz")
+
+npzOCLLG2=numpy.load("Ising2D_OCLL_256_16777216_GPU_o745-2.npz")
+npzCUDAL2=numpy.load("Ising2D_CUDAL_256_16777216_GPU_o745-2.npz")
+
+
+T=npzOCLLG['arr_0'][0]
+EOCLLG=npzOCLLG['arr_0'][1]
+POCLLG=npzOCLLG['arr_0'][2]
+EOCLLC=npzOCLLC['arr_0'][1]
+POCLLC=npzOCLLC['arr_0'][2]
+ECUDAL=npzCUDAL['arr_0'][1]
+PCUDAL=npzCUDAL['arr_0'][2]
+
+EOCLLG2=npzOCLLG2['arr_0'][1]
+POCLLG2=npzOCLLG2['arr_0'][2]
+ECUDAL2=npzCUDAL2['arr_0'][1]
+PCUDAL2=npzCUDAL2['arr_0'][2]
+
+dPOCLLG=numpy.diff(POCLLG)
+dPOCLLC=numpy.diff(POCLLC)
+dPCUDAL=numpy.diff(PCUDAL)
+
+dPOCLLG2=numpy.diff(POCLLG2)
+dPCUDAL2=numpy.diff(PCUDAL2)
+
+dPOCLLG=numpy.insert(dPOCLLG,0,0)
+dPOCLLC=numpy.insert(dPOCLLC,0,0)
+dPCUDAL=numpy.insert(dPCUDAL,0,0)
+
+dPOCLLG2=numpy.insert(dPOCLLG2,0,0)
+dPCUDAL2=numpy.insert(dPCUDAL2,0,0)
+
+from matplotlib.pyplot import *
+
+plot(T,EOCLLG,T,EOCLLC,T,ECUDAL,T,dPOCLLG,T,dPOCLLC,T,dPCUDAL,
+     T,EOCLLG2,T,ECUDAL2,T,dPOCLLG2,T,dPCUDAL2)
+legend((r'OCL GPU GTX560',r'OCL CPU Xeon4c',r'Cuda GTX560',
+        r'OCL GPU GT8400',r'Cuda GT8400'))
+xlabel('Temperature',{'fontsize':20})
+ylabel('Energy',{'fontsize':20})
+
+matplotlib.pyplot.show()
+
Index: Pi/Results/PostProcess.py
===================================================================
--- Pi/Results/PostProcess.py	(revision 0)
+++ Pi/Results/PostProcess.py	(revision 7)
@@ -0,0 +1,24 @@
+import numpy
+import sys
+
+npzSerial=numpy.load("Ising2D_Serial_256_16777216.npz")
+npzMP=numpy.load("Ising2D_MP_256_16777216.npz")
+npzMPI=numpy.load("Ising2D_MPI_256_16777216.npz")
+npzOCLG=numpy.load("Ising2D_OCLG_256_16777216.npz")
+npzOCLL=numpy.load("Ising2D_OCLL_256_16777216.npz")
+npzCUDAG=numpy.load("Ising2D_CUDAG_256_16777216.npz")
+npzCUDAL=numpy.load("Ising2D_CUDAL_256_16777216.npz")
+
+T=npzSerial['arr_0'][0]
+ES=npzSerial['arr_0'][1]
+EMP=npzMP['arr_0'].transpose()[1]
+EMPI=npzMPI['arr_0'].transpose()[1]
+EOCLG=npzOCLG['arr_0'][1]
+EOCLL=npzOCLL['arr_0'][1]
+ECUDAG=npzCUDAG['arr_0'][1]
+ECUDAL=npzCUDAL['arr_0'][1]
+
+import matplotlib.pyplot
+matplotlib.pyplot.plot(T,ES,T,EMP,T,EMPI,T,EOCLG,T,EOCLL,T,ECUDAG,T,ECUDAL)
+matplotlib.pyplot.show()
+
Index: Pi/Results/PostProcessGPU.py
===================================================================
--- Pi/Results/PostProcessGPU.py	(revision 0)
+++ Pi/Results/PostProcessGPU.py	(revision 7)
@@ -0,0 +1,20 @@
+import numpy
+import sys
+
+npzSerial=numpy.load("Ising2D_Serial_256_16777216.npz")
+npzOCLG=numpy.load("Ising2D_OCLG_256_16777216.npz")
+npzOCLL=numpy.load("Ising2D_OCLL_256_16777216.npz")
+npzCUDAG=numpy.load("Ising2D_CUDAG_256_16777216.npz")
+npzCUDAL=numpy.load("Ising2D_CUDAL_256_16777216.npz")
+
+T=npzSerial['arr_0'][0]
+ES=npzSerial['arr_0'][1]
+EOCLG=npzOCLG['arr_0'][1]
+EOCLL=npzOCLL['arr_0'][1]
+ECUDAG=npzCUDAG['arr_0'][1]
+ECUDAL=npzCUDAL['arr_0'][1]
+
+import matplotlib.pyplot
+matplotlib.pyplot.plot(T,ES,T,EOCLG,T,EOCLL,T,ECUDAG,T,ECUDAL)
+matplotlib.pyplot.show()
+
Index: Pi/Results/PostProcessMP.py
===================================================================
--- Pi/Results/PostProcessMP.py	(revision 0)
+++ Pi/Results/PostProcessMP.py	(revision 7)
@@ -0,0 +1,16 @@
+import numpy
+import sys
+
+npzSerial=numpy.load("Ising2D_Serial_256_16777216.npz")
+npzMP=numpy.load("Ising2D_MP_256_16777216.npz")
+npzMPI=numpy.load("Ising2D_MPI_256_16777216.npz")
+
+T=npzSerial['arr_0'][0]
+ES=npzSerial['arr_0'][1]
+EMP=npzMP['arr_0'].transpose()[1]
+EMPI=npzMPI['arr_0'].transpose()[1]
+
+import matplotlib.pyplot
+matplotlib.pyplot.plot(T,ES,T,EMP,T,EMPI)
+matplotlib.pyplot.show()
+
Index: Pi/C/SWIG/MC.py
===================================================================
--- Pi/C/SWIG/MC.py	(revision 0)
+++ Pi/C/SWIG/MC.py	(revision 7)
@@ -0,0 +1,7 @@
+#!/usr/bin/env python
+
+import MonteCarlo
+
+print MonteCarlo.InsideCircle(1000000,1,2)
+
+
Index: Pi/C/SWIG/MCmodule.c
===================================================================
--- Pi/C/SWIG/MCmodule.c	(revision 0)
+++ Pi/C/SWIG/MCmodule.c	(revision 7)
@@ -0,0 +1,35 @@
+#include <Python.h>
+#include <math.h>
+
+// 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
+
+int InsideCircle(unsigned int iterations,
+		 unsigned int seed_w,unsigned int seed_z)
+{
+   unsigned int z=seed_z;
+   unsigned int w=seed_w;
+   unsigned int i;
+
+   int total=0;
+
+   for (i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   return(total);
+}
Index: Pi/C/SWIG/Build.sh
===================================================================
--- Pi/C/SWIG/Build.sh	(revision 0)
+++ Pi/C/SWIG/Build.sh	(revision 7)
@@ -0,0 +1,12 @@
+#!/bin/bash
+# Delete files created by Swig compilation process
+rm _MonteCarlo.so MonteCarlo_wrap.o MonteCarlo_wrap.c MonteCarlo.py MonteCarlo.pyc MCmodule.o
+# MonteCarlo.i holds prototype (like .h) of InsideCircle function
+# Swig creates MonteCarlo_wrap.c and MonteCarlo.py
+swig -python MonteCarlo.i
+# Compilation of MCmodule.c as wrapper
+gcc -O3 -fpic -c MonteCarlo_wrap.c MCmodule.c -I/usr/include/python2.7/
+# Link to create library
+gcc -O3 -shared MCmodule.o MonteCarlo_wrap.o -o _MonteCarlo.so
+# Execution
+python MC.py
\ No newline at end of file

Property changes on: Pi/C/SWIG/Build.sh
___________________________________________________________________
Added: svn:executable
   + *

Index: Pi/C/SWIG/MonteCarlo.i
===================================================================
--- Pi/C/SWIG/MonteCarlo.i	(revision 0)
+++ Pi/C/SWIG/MonteCarlo.i	(revision 7)
@@ -0,0 +1,2 @@
+%module MonteCarlo
+extern int InsideCircle(unsigned int iterations,unsigned int seed_w,unsigned int seed_z);
Index: Pi/C/Simple/Pi_INTEGER
===================================================================
Cannot display: file marked as a binary type.
svn:mime-type = application/octet-stream

Property changes on: Pi/C/Simple/Pi_INTEGER
___________________________________________________________________
Added: svn:mime-type
   + application/octet-stream
Added: svn:executable
   + *

Index: Pi/C/Simple/Pi_Integer.c
===================================================================
--- Pi/C/Simple/Pi_Integer.c	(revision 0)
+++ Pi/C/Simple/Pi_Integer.c	(revision 7)
@@ -0,0 +1,59 @@
+//
+// Estimation of Pi using Monte Carlo exploration process
+// gcc -std=c99 -O3 -o Pi Pi.c -lm 
+//
+#include <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+// 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 ITERATIONS 1000000000
+
+unsigned int MainLoopGlobal(unsigned int iterations,unsigned int seed_w,unsigned int seed_z)
+{
+   unsigned int z=seed_z;
+   unsigned int w=seed_w;
+
+   unsigned int total=0;
+
+   for (unsigned int i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   return(total);
+
+}
+
+int main(int argc, char *argv[]) {
+
+  unsigned int seed_w=10,seed_z=10;
+  unsigned int iterations=ITERATIONS;
+
+  if (argc > 1) {
+    iterations=(unsigned int)atol(argv[1]);
+  }
+  else {
+    printf("\n\tPi : Estimate Pi with Monte Carlo exploration\n\n\t\t#1 : number of iterations (default 1 billion)\n");
+  }
+
+  float pi=(float)MainLoopGlobal(iterations,seed_w,seed_z)/(float)iterations*4;
+
+  printf("Pi=%f with %u iterations\n",pi,iterations);
+  
+}
Index: Pi/C/Simple/Pi_LONG
===================================================================
Cannot display: file marked as a binary type.
svn:mime-type = application/octet-stream

Property changes on: Pi/C/Simple/Pi_LONG
___________________________________________________________________
Added: svn:mime-type
   + application/octet-stream
Added: svn:executable
   + *

Index: Pi/C/Simple/Pi_Long.c
===================================================================
--- Pi/C/Simple/Pi_Long.c	(revision 0)
+++ Pi/C/Simple/Pi_Long.c	(revision 7)
@@ -0,0 +1,59 @@
+//
+// Estimation of Pi using Monte Carlo exploration process
+// gcc -std=c99 -O3 -o Pi Pi.c -lm 
+//
+#include <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+// 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 ITERATIONS 1000000000
+
+unsigned long MainLoopGlobal(unsigned long iterations,unsigned int seed_w,unsigned int seed_z)
+{
+   unsigned int z=seed_z;
+   unsigned int w=seed_w;
+
+   unsigned  long total=0;
+
+   for (unsigned long i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   return(total);
+
+}
+
+int main(int argc, char *argv[]) {
+
+  unsigned int seed_w=10,seed_z=10;
+  unsigned long iterations=ITERATIONS;
+
+  if (argc > 1) {
+    iterations=(unsigned long)atol(argv[1]);
+  }
+  else {
+    printf("\n\tPi : Estimate Pi with Monte Carlo exploration\n\n\t\t#1 : number of iterations (default 1 billion)\n");
+  }
+
+  float pi=(float)MainLoopGlobal(iterations,seed_w,seed_z)/(float)iterations*4;
+
+  printf("Pi=%f with %lu iterations\n",pi,iterations);
+  
+}
Index: Pi/C/Simple/Pi.c
===================================================================
--- Pi/C/Simple/Pi.c	(revision 0)
+++ Pi/C/Simple/Pi.c	(revision 7)
@@ -0,0 +1,72 @@
+//
+// Estimation of Pi using Monte Carlo exploration process
+// gcc -std=c99 -O3 -o Pi Pi.c -lm 
+//
+
+#include <math.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+// 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 ITERATIONS 1000000000
+
+#ifdef LONG
+#define LENGTH unsigned long
+#else
+#define LENGTH unsigned int
+#endif
+
+LENGTH MainLoopGlobal(LENGTH iterations,unsigned int seed_w,unsigned int seed_z)
+{
+   unsigned int z=seed_z;
+   unsigned int w=seed_w;
+
+   unsigned  long total=0;
+
+   for (LENGTH i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   return(total);
+
+}
+
+int main(int argc, char *argv[]) {
+
+  unsigned int seed_w=10,seed_z=10;
+  LENGTH iterations=ITERATIONS;
+
+  if (argc > 1) {
+    iterations=(LENGTH)atol(argv[1]);
+  }
+  else {
+    printf("\n\tPi : Estimate Pi with Monte Carlo exploration\n\n\t\t#1 : number of iterations (default 1 billion)\n\n");
+  }
+
+  float pi=(float)MainLoopGlobal(iterations,seed_w,seed_z)/(float)iterations*4;
+
+#ifdef LONG
+  printf("\tPi=%f with error %f and %lu iterations\n\n",pi,
+         fabs(pi-4*atan(1))/pi,iterations);
+#else
+  printf("\tPi=%f with error %f and %u iterations\n\n",pi,
+         fabs(pi-4*atan(1))/pi,iterations);
+#endif
+  
+}
Index: Pi/C/Simple/Makefile
===================================================================
--- Pi/C/Simple/Makefile	(revision 0)
+++ Pi/C/Simple/Makefile	(revision 7)
@@ -0,0 +1,30 @@
+EXECUTABLE=Pi_LONG Pi_INTEGER
+
+SOURCE=Pi.c
+
+CC=gcc
+CFLAGS=-Wall -O3 -std=c99
+LIBRARY=-lm
+
+all: $(EXECUTABLE)
+
+Pi_LONG: $(SOURCE)
+
+	$(CC) $(CFLAGS) $(DIRECTIVES) -DLONG -o $@ $< $(LIBRARY)
+
+Pi_INTEGER: $(SOURCE)
+
+	$(CC) $(CFLAGS) $(DIRECTIVES) -DINTEGER -o $@ $< $(LIBRARY)
+
+.PHONY: clean check mrproper
+
+mrproper: 
+	rm -rf $(EXECUTABLE)
+	find . -name "*~" -exec rm {} \;
+
+clean:
+	find . -name "*~" -exec rm {} \;
+
+check: $(EXECUTABLE)
+
+	$(SOURCE:.c=)_SP_$(EXECUTABLE) 2 1
Index: Pi/C/Pthreads/PiMP.c
===================================================================
--- Pi/C/Pthreads/PiMP.c	(revision 0)
+++ Pi/C/Pthreads/PiMP.c	(revision 7)
@@ -0,0 +1,155 @@
+/******************************************************************************
+* FILE: hello.c
+* DESCRIPTION:
+*   A "hello world" Pthreads program.  Demonstrates thread creation and
+*   termination.
+* AUTHOR: Blaise Barney
+* LAST REVISED: 08/09/11
+******************************************************************************/
+#include <pthread.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <stdio.h>
+
+#define NUM_THREADS	1024
+
+struct thread_data
+{
+  int thread_id;
+  unsigned int seed_w;
+  unsigned int seed_z;
+  unsigned long iterations;
+  unsigned long inside;
+};
+
+struct thread_data thread_data_array[NUM_THREADS];
+
+// 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 ITERATIONS 1000000000
+
+int MainLoopGlobal(unsigned long iterations,unsigned int seed_w,unsigned int seed_z)
+{
+  unsigned int z=seed_z;
+  unsigned int w=seed_w;
+  
+  unsigned long total=0;
+  
+  for (unsigned long i=0;i<iterations;i++) {
+    
+    float x=MWCfp ;
+    float y=MWCfp ;
+    
+    // Matching test
+    unsigned long inside=((x*x+y*y) < 1.0f) ? 1:0;
+    total+=inside;
+   }
+  
+  return(total);
+  
+}
+
+void *MainLoopThread(void *threadarg)
+{
+  int taskid;
+  unsigned long iterations,total=0;
+  unsigned int z,w;
+  
+  struct thread_data *my_data;
+  
+  my_data=(struct thread_data *) threadarg;
+  
+  taskid = my_data->thread_id;
+  iterations = my_data->iterations;
+  z = my_data->seed_z;
+  w = my_data->seed_w;
+  
+  printf("Hello World! It's me, thread #%ld, with seeds (%ld,%ld) and %ld !\n", taskid,z,w,iterations);
+  
+  for (unsigned long i=0;i<iterations;i++) {
+    
+    float x=MWCfp ;
+    float y=MWCfp ;
+    
+    // Matching test
+    unsigned long inside=((x*x+y*y) < 1.0f) ? 1:0;
+    total+=inside;
+  }
+  
+  my_data->inside=total;
+  
+  printf("Thread %ld done, found %ld inside.\n",taskid,total);
+  
+  pthread_exit((void*) my_data);
+}
+
+int main(int argc, char *argv[])
+{
+  pthread_t threads[NUM_THREADS];
+  pthread_attr_t attr;
+  int *taskids[NUM_THREADS];
+  int rc, t, sum;
+  unsigned long num_threads,iterations;
+  sum=0;
+  void *status;
+  
+  iterations=atol(argv[1]);
+  num_threads=atoi(argv[2]);
+  
+  printf("Number of threads defined to %ld\n",num_threads);
+  printf("Number of iterations defined to %ld\n",iterations);
+
+  /* Initialize and set thread detached attribute */
+  pthread_attr_init(&attr);
+  pthread_attr_setdetachstate(&attr, PTHREAD_CREATE_JOINABLE);
+
+  for(t=0;t<num_threads;t++) {
+    
+    thread_data_array[t].thread_id = t;
+    thread_data_array[t].iterations = iterations/num_threads;
+    thread_data_array[t].seed_w = (t+1)<<4;
+    thread_data_array[t].seed_z = (1048576*(t+1))>>4;
+    thread_data_array[t].inside = 0;
+    
+    printf("Creating thread %d\n", t);
+    rc = pthread_create(&threads[t], &attr, MainLoopThread, (void *) 
+                        &thread_data_array[t]);
+
+    if (rc) {
+      printf("ERROR; return code from pthread_create() is %d\n", rc);
+      exit(-1);
+    }
+  }
+  
+  /* Free attribute and wait for the other threads */
+  pthread_attr_destroy(&attr);
+
+  for(t=0; t<num_threads; t++) {
+    rc = pthread_join(threads[t], &status);
+    if (rc) {
+      printf("ERROR; return code from pthread_join() is %d\n", rc);
+      exit(-1);
+    }
+    printf("Main: completed join with thread %ld having a status of %ld\n",
+           t,(long)status);
+  }
+
+  for(t=0;t<num_threads;t++) {
+    printf("Return to main with %ld thread, %ld inside\n", 
+           t,thread_data_array[t].inside);
+  }
+
+  printf("Main: program completed. Exiting.\n");
+  pthread_exit(NULL);
+
+}
Index: Pi/C/Pthreads/PiMP_orig1.c
===================================================================
--- Pi/C/Pthreads/PiMP_orig1.c	(revision 0)
+++ Pi/C/Pthreads/PiMP_orig1.c	(revision 7)
@@ -0,0 +1,104 @@
+/******************************************************************************
+* FILE: hello.c
+* DESCRIPTION:
+*   A "hello world" Pthreads program.  Demonstrates thread creation and
+*   termination.
+* AUTHOR: Blaise Barney
+* LAST REVISED: 08/09/11
+******************************************************************************/
+#include <pthread.h>
+#include <stdio.h>
+#include <stdlib.h>
+#define NUM_THREADS	4
+
+#include <math.h>
+#include <stdio.h>
+
+// 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 ITERATIONS 1000000000
+
+int MainLoopGlobal(unsigned int iterations,unsigned int seed_w,unsigned int seed_z)
+{
+   unsigned int z=seed_z;
+   unsigned int w=seed_w;
+
+   int total=0;
+
+   for (unsigned int i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   return(total);
+
+}
+
+void *PrintHello(void *threadid)
+{
+   long tid;
+   tid = (long)threadid;
+   printf("Hello World! It's me, thread #%ld!\n", tid);
+
+   pthread_exit(NULL);
+}
+
+void *MainLoopThread(void *threadid)
+{
+   long tid;
+   tid = (long)threadid;
+   printf("Hello World! It's me, thread #%ld!\n", tid);
+
+   /* unsigned int z=seed_z; */
+   /* unsigned int w=seed_w; */
+   unsigned int z=16;
+   unsigned int w=64;
+
+   int total=0;
+
+   for (unsigned int i=0;i<ITERATIONS;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   printf("In #%ld, found %i inside\n",tid,total);
+   
+   pthread_exit(NULL);
+}
+
+int main(int argc, char *argv[])
+{
+   pthread_t threads[NUM_THREADS];
+   int rc;
+   long t;
+   for(t=0;t<NUM_THREADS;t++){
+     printf("In main: creating thread %ld\n", t);
+     rc = pthread_create(&threads[t], NULL, MainLoopThread, (void *)t);
+     if (rc){
+       printf("ERROR; return code from pthread_create() is %d\n", rc);
+       exit(-1);
+       }
+     }
+
+   /* Last thing that main() should do */
+   pthread_exit(NULL);
+}
Index: Pi/C/Pthreads/PiMP_orig2.c
===================================================================
--- Pi/C/Pthreads/PiMP_orig2.c	(revision 0)
+++ Pi/C/Pthreads/PiMP_orig2.c	(revision 7)
@@ -0,0 +1,125 @@
+/******************************************************************************
+* FILE: hello.c
+* DESCRIPTION:
+*   A "hello world" Pthreads program.  Demonstrates thread creation and
+*   termination.
+* AUTHOR: Blaise Barney
+* LAST REVISED: 08/09/11
+******************************************************************************/
+#include <pthread.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <stdio.h>
+
+#define NUM_THREADS	4
+
+struct thread_data
+{
+  int thread_id;
+  unsigned int seed_w;
+  unsigned int seed_z;
+  unsigned int iterations;
+  unsigned int inside;
+};
+
+struct thread_data thread_data_array[NUM_THREADS];
+
+// 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 ITERATIONS 1000000000
+
+int MainLoopGlobal(unsigned int iterations,unsigned int seed_w,unsigned int seed_z)
+{
+   unsigned int z=seed_z;
+   unsigned int w=seed_w;
+
+   int total=0;
+
+   for (unsigned int i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   return(total);
+
+}
+
+void *MainLoopThread(void *threadarg)
+{
+  int taskid;
+  unsigned int iterations,z,w,total=0;
+
+  struct thread_data *my_data;
+
+  my_data=(struct thread_data *) threadarg;
+
+  taskid = my_data->thread_id;
+  iterations = my_data->iterations;
+  z = my_data->seed_z;
+  w = my_data->seed_w;
+  iterations = my_data->iterations;
+
+  printf("Hello World! It's me, thread #%ld, with seeds (%i,%i) and %i !\n", taskid,z,w,iterations);
+
+   for (unsigned int i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   printf("In #%ld, found %i inside\n",taskid,total);
+   
+   my_data->inside=total;
+
+   pthread_exit(NULL);
+}
+
+int main(int argc, char *argv[])
+{
+
+pthread_t threads[NUM_THREADS];
+int *taskids[NUM_THREADS];
+int rc, t, sum;
+
+sum=0;
+
+for(t=0;t<NUM_THREADS;t++) {
+
+  thread_data_array[t].thread_id = t;
+  thread_data_array[t].iterations = ITERATIONS/NUM_THREADS;
+  thread_data_array[t].seed_w = (t+1)<<4;
+  thread_data_array[t].seed_z = (1048576*(t+1))>>4;
+  thread_data_array[t].inside = 0;
+
+  printf("Creating thread %d\n", t);
+  rc = pthread_create(&threads[t], NULL, MainLoopThread, (void *) 
+       &thread_data_array[t]);
+  printf("%i inside\n", thread_data_array[t].inside);
+  if (rc) {
+    printf("ERROR; return code from pthread_create() is %d\n", rc);
+    exit(-1);
+    }
+  }
+pthread_exit(NULL);
+
+
+}
Index: Pi/C/Pthreads/PiMP_orig3.c
===================================================================
--- Pi/C/Pthreads/PiMP_orig3.c	(revision 0)
+++ Pi/C/Pthreads/PiMP_orig3.c	(revision 7)
@@ -0,0 +1,128 @@
+/******************************************************************************
+* FILE: hello.c
+* DESCRIPTION:
+*   A "hello world" Pthreads program.  Demonstrates thread creation and
+*   termination.
+* AUTHOR: Blaise Barney
+* LAST REVISED: 08/09/11
+******************************************************************************/
+#include <pthread.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <math.h>
+#include <stdio.h>
+
+#define NUM_THREADS	1024
+
+struct thread_data
+{
+  int thread_id;
+  unsigned int seed_w;
+  unsigned int seed_z;
+  unsigned int iterations;
+  unsigned int inside;
+};
+
+struct thread_data thread_data_array[NUM_THREADS];
+
+// 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 ITERATIONS 1000000000
+
+int MainLoopGlobal(unsigned int iterations,unsigned int seed_w,unsigned int seed_z)
+{
+   unsigned int z=seed_z;
+   unsigned int w=seed_w;
+
+   int total=0;
+
+   for (unsigned int i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   return(total);
+
+}
+
+void *MainLoopThread(void *threadarg)
+{
+  int taskid;
+  unsigned int iterations,z,w,total=0;
+
+  struct thread_data *my_data;
+
+  my_data=(struct thread_data *) threadarg;
+
+  taskid = my_data->thread_id;
+  iterations = my_data->iterations;
+  z = my_data->seed_z;
+  w = my_data->seed_w;
+  iterations = my_data->iterations;
+
+  printf("Hello World! It's me, thread #%ld, with seeds (%i,%i) and %i !\n", taskid,z,w,iterations);
+
+   for (unsigned int i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   printf("In #%ld, found %i inside\n",taskid,total);
+   
+   my_data->inside=total;
+
+   pthread_exit(NULL);
+}
+
+int main(int argc, char *argv[])
+{
+
+  pthread_t threads[NUM_THREADS];
+  int *taskids[NUM_THREADS];
+  int rc, t, sum;
+  int num_threads,iterations;
+  sum=0;
+
+  iterations=atoi(argv[1]);
+  num_threads=atoi(argv[2]);
+
+  printf("Number of threads defined to %i\n",num_threads);
+
+for(t=0;t<num_threads;t++) {
+  
+  thread_data_array[t].thread_id = t;
+  thread_data_array[t].iterations = iterations/num_threads;
+  thread_data_array[t].seed_w = (t+1)<<4;
+  thread_data_array[t].seed_z = (1048576*(t+1))>>4;
+  thread_data_array[t].inside = 0;
+  
+  printf("Creating thread %d\n", t);
+  rc = pthread_create(&threads[t], NULL, MainLoopThread, (void *) 
+                      &thread_data_array[t]);
+  printf("%i inside\n", thread_data_array[t].inside);
+  if (rc) {
+    printf("ERROR; return code from pthread_create() is %d\n", rc);
+    exit(-1);
+  }
+ }
+ pthread_exit(NULL);
+}
Index: Pi/C/CAPI/pimodule.c
===================================================================
--- Pi/C/CAPI/pimodule.c	(revision 0)
+++ Pi/C/CAPI/pimodule.c	(revision 7)
@@ -0,0 +1,56 @@
+#include <Python.h>
+#include <math.h>
+
+// 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
+
+int _InsideCircle(unsigned int iterations,
+		  unsigned int seed_w,unsigned int seed_z)
+{
+   unsigned int z=seed_z;
+   unsigned int w=seed_w;
+   unsigned int i;
+
+   int total=0;
+
+   for (i=0;i<iterations;i++) {
+
+      float x=MWCfp ;
+      float y=MWCfp ;
+
+      // Matching test
+      int inside=((x*x+y*y) < 1.0f) ? 1:0;
+      total+=inside;
+   }
+
+   return(total);
+}
+
+static PyObject* InsideCircle(PyObject* self, PyObject* args)
+{
+  int i,z,w;
+ 
+  if (!PyArg_ParseTuple(args, "iii", &i,&z,&w))
+    return NULL;
+ 
+  return Py_BuildValue("i", _InsideCircle(i,z,w));
+}
+ 
+static PyMethodDef MonteCarloMethods[] = {
+  {"InsideCircle", InsideCircle, METH_VARARGS, "Inside Circle Randomly Check"},
+  {NULL, NULL, 0, NULL}
+};
+ 
+PyMODINIT_FUNC
+initMonteCarlo(void)
+{
+  (void) Py_InitModule("MonteCarlo", MonteCarloMethods);
+}
Index: Pi/C/CAPI/MC.py
===================================================================
--- Pi/C/CAPI/MC.py	(revision 0)
+++ Pi/C/CAPI/MC.py	(revision 7)
@@ -0,0 +1,5 @@
+#!/usr/bin/env python
+
+import MonteCarlo
+
+print MonteCarlo.InsideCircle(1000000000,1,2)
Index: Pi/C/CAPI/setup.py
===================================================================
--- Pi/C/CAPI/setup.py	(revision 0)
+++ Pi/C/CAPI/setup.py	(revision 7)
@@ -0,0 +1,8 @@
+from distutils.core import setup, Extension
+
+module1 = Extension('MonteCarlo', sources = ['pimodule.c'])
+
+setup (name = 'PackageName',
+       version = '1.0',
+       description = 'This is a demo package',
+       ext_modules = [module1])
Index: Pi/C/CAPI/Build.sh
===================================================================
--- Pi/C/CAPI/Build.sh	(revision 0)
+++ Pi/C/CAPI/Build.sh	(revision 7)
@@ -0,0 +1,4 @@
+rm -r build MonteCarlo.so 
+python setup.py build
+cp build/lib.linux-x86_64-2.7/MonteCarlo.so .
+python MC.py

Property changes on: Pi/C/CAPI/Build.sh
___________________________________________________________________
Added: svn:executable
   + *

Index: Pi/C/examples/swig/hello.i
===================================================================
--- Pi/C/examples/swig/hello.i	(revision 0)
+++ Pi/C/examples/swig/hello.i	(revision 7)
@@ -0,0 +1,4 @@
+/*hello.i*/
+ 
+%module hello
+extern void say_hello(const char* name);
Index: Pi/C/examples/swig/hello.py
===================================================================
--- Pi/C/examples/swig/hello.py	(revision 0)
+++ Pi/C/examples/swig/hello.py	(revision 7)
@@ -0,0 +1,76 @@
+# This file was automatically generated by SWIG (http://www.swig.org).
+# Version 2.0.7
+#
+# Do not make changes to this file unless you know what you are doing--modify
+# the SWIG interface file instead.
+
+
+
+from sys import version_info
+if version_info >= (2,6,0):
+    def swig_import_helper():
+        from os.path import dirname
+        import imp
+        fp = None
+        try:
+            fp, pathname, description = imp.find_module('_hello', [dirname(__file__)])
+        except ImportError:
+            import _hello
+            return _hello
+        if fp is not None:
+            try:
+                _mod = imp.load_module('_hello', fp, pathname, description)
+            finally:
+                fp.close()
+            return _mod
+    _hello = swig_import_helper()
+    del swig_import_helper
+else:
+    import _hello
+del version_info
+try:
+    _swig_property = property
+except NameError:
+    pass # Python < 2.2 doesn't have 'property'.
+def _swig_setattr_nondynamic(self,class_type,name,value,static=1):
+    if (name == "thisown"): return self.this.own(value)
+    if (name == "this"):
+        if type(value).__name__ == 'SwigPyObject':
+            self.__dict__[name] = value
+            return
+    method = class_type.__swig_setmethods__.get(name,None)
+    if method: return method(self,value)
+    if (not static):
+        self.__dict__[name] = value
+    else:
+        raise AttributeError("You cannot add attributes to %s" % self)
+
+def _swig_setattr(self,class_type,name,value):
+    return _swig_setattr_nondynamic(self,class_type,name,value,0)
+
+def _swig_getattr(self,class_type,name):
+    if (name == "thisown"): return self.this.own()
+    method = class_type.__swig_getmethods__.get(name,None)
+    if method: return method(self)
+    raise AttributeError(name)
+
+def _swig_repr(self):
+    try: strthis = "proxy of " + self.this.__repr__()
+    except: strthis = ""
+    return "<%s.%s; %s >" % (self.__class__.__module__, self.__class__.__name__, strthis,)
+
+try:
+    _object = object
+    _newclass = 1
+except AttributeError:
+    class _object : pass
+    _newclass = 0
+
+
+
+def say_hello(*args):
+  return _hello.say_hello(*args)
+say_hello = _hello.say_hello
+# This file is compatible with both classic and new-style classes.
+
+
Index: Pi/C/examples/swig/hellomodule.c
===================================================================
--- Pi/C/examples/swig/hellomodule.c	(revision 0)
+++ Pi/C/examples/swig/hellomodule.c	(revision 7)
@@ -0,0 +1,7 @@
+/*hellomodule.c*/
+ 
+#include <stdio.h>
+ 
+void say_hello(const char* name) {
+    printf("Hello %s!\n", name);
+}
Index: Pi/C/examples/swig/hello_wrap.c
===================================================================
--- Pi/C/examples/swig/hello_wrap.c	(revision 0)
+++ Pi/C/examples/swig/hello_wrap.c	(revision 7)
@@ -0,0 +1,3793 @@
+/* ----------------------------------------------------------------------------
+ * This file was automatically generated by SWIG (http://www.swig.org).
+ * Version 2.0.7
+ * 
+ * This file is not intended to be easily readable and contains a number of 
+ * coding conventions designed to improve portability and efficiency. Do not make
+ * changes to this file unless you know what you are doing--modify the SWIG 
+ * interface file instead. 
+ * ----------------------------------------------------------------------------- */
+
+#define SWIGPYTHON
+#define SWIG_PYTHON_DIRECTOR_NO_VTABLE
+
+/* -----------------------------------------------------------------------------
+ *  This section contains generic SWIG labels for method/variable
+ *  declarations/attributes, and other compiler dependent labels.
+ * ----------------------------------------------------------------------------- */
+
+/* template workaround for compilers that cannot correctly implement the C++ standard */
+#ifndef SWIGTEMPLATEDISAMBIGUATOR
+# if defined(__SUNPRO_CC) && (__SUNPRO_CC <= 0x560)
+#  define SWIGTEMPLATEDISAMBIGUATOR template
+# elif defined(__HP_aCC)
+/* Needed even with `aCC -AA' when `aCC -V' reports HP ANSI C++ B3910B A.03.55 */
+/* If we find a maximum version that requires this, the test would be __HP_aCC <= 35500 for A.03.55 */
+#  define SWIGTEMPLATEDISAMBIGUATOR template
+# else
+#  define SWIGTEMPLATEDISAMBIGUATOR
+# endif
+#endif
+
+/* inline attribute */
+#ifndef SWIGINLINE
+# if defined(__cplusplus) || (defined(__GNUC__) && !defined(__STRICT_ANSI__))
+#   define SWIGINLINE inline
+# else
+#   define SWIGINLINE
+# endif
+#endif
+
+/* attribute recognised by some compilers to avoid 'unused' warnings */
+#ifndef SWIGUNUSED
+# if defined(__GNUC__)
+#   if !(defined(__cplusplus)) || (__GNUC__ > 3 || (__GNUC__ == 3 && __GNUC_MINOR__ >= 4))
+#     define SWIGUNUSED __attribute__ ((__unused__)) 
+#   else
+#     define SWIGUNUSED
+#   endif
+# elif defined(__ICC)
+#   define SWIGUNUSED __attribute__ ((__unused__)) 
+# else
+#   define SWIGUNUSED 
+# endif
+#endif
+
+#ifndef SWIG_MSC_UNSUPPRESS_4505
+# if defined(_MSC_VER)
+#   pragma warning(disable : 4505) /* unreferenced local function has been removed */
+# endif 
+#endif
+
+#ifndef SWIGUNUSEDPARM
+# ifdef __cplusplus
+#   define SWIGUNUSEDPARM(p)
+# else
+#   define SWIGUNUSEDPARM(p) p SWIGUNUSED 
+# endif
+#endif
+
+/* internal SWIG method */
+#ifndef SWIGINTERN
+# define SWIGINTERN static SWIGUNUSED
+#endif
+
+/* internal inline SWIG method */
+#ifndef SWIGINTERNINLINE
+# define SWIGINTERNINLINE SWIGINTERN SWIGINLINE
+#endif
+
+/* exporting methods */
+#if (__GNUC__ >= 4) || (__GNUC__ == 3 && __GNUC_MINOR__ >= 4)
+#  ifndef GCC_HASCLASSVISIBILITY
+#    define GCC_HASCLASSVISIBILITY
+#  endif
+#endif
+
+#ifndef SWIGEXPORT
+# if defined(_WIN32) || defined(__WIN32__) || defined(__CYGWIN__)
+#   if defined(STATIC_LINKED)
+#     define SWIGEXPORT
+#   else
+#     define SWIGEXPORT __declspec(dllexport)
+#   endif
+# else
+#   if defined(__GNUC__) && defined(GCC_HASCLASSVISIBILITY)
+#     define SWIGEXPORT __attribute__ ((visibility("default")))
+#   else
+#     define SWIGEXPORT
+#   endif
+# endif
+#endif
+
+/* calling conventions for Windows */
+#ifndef SWIGSTDCALL
+# if defined(_WIN32) || defined(__WIN32__) || defined(__CYGWIN__)
+#   define SWIGSTDCALL __stdcall
+# else
+#   define SWIGSTDCALL
+# endif 
+#endif
+
+/* Deal with Microsoft's attempt at deprecating C standard runtime functions */
+#if !defined(SWIG_NO_CRT_SECURE_NO_DEPRECATE) && defined(_MSC_VER) && !defined(_CRT_SECURE_NO_DEPRECATE)
+# define _CRT_SECURE_NO_DEPRECATE
+#endif
+
+/* Deal with Microsoft's attempt at deprecating methods in the standard C++ library */
+#if !defined(SWIG_NO_SCL_SECURE_NO_DEPRECATE) && defined(_MSC_VER) && !defined(_SCL_SECURE_NO_DEPRECATE)
+# define _SCL_SECURE_NO_DEPRECATE
+#endif
+
+
+
+/* Python.h has to appear first */
+#include <Python.h>
+
+/* -----------------------------------------------------------------------------
+ * swigrun.swg
+ *
+ * This file contains generic C API SWIG runtime support for pointer
+ * type checking.
+ * ----------------------------------------------------------------------------- */
+
+/* This should only be incremented when either the layout of swig_type_info changes,
+   or for whatever reason, the runtime changes incompatibly */
+#define SWIG_RUNTIME_VERSION "4"
+
+/* define SWIG_TYPE_TABLE_NAME as "SWIG_TYPE_TABLE" */
+#ifdef SWIG_TYPE_TABLE
+# define SWIG_QUOTE_STRING(x) #x
+# define SWIG_EXPAND_AND_QUOTE_STRING(x) SWIG_QUOTE_STRING(x)
+# define SWIG_TYPE_TABLE_NAME SWIG_EXPAND_AND_QUOTE_STRING(SWIG_TYPE_TABLE)
+#else
+# define SWIG_TYPE_TABLE_NAME
+#endif
+
+/*
+  You can use the SWIGRUNTIME and SWIGRUNTIMEINLINE macros for
+  creating a static or dynamic library from the SWIG runtime code.
+  In 99.9% of the cases, SWIG just needs to declare them as 'static'.
+  
+  But only do this if strictly necessary, ie, if you have problems
+  with your compiler or suchlike.
+*/
+
+#ifndef SWIGRUNTIME
+# define SWIGRUNTIME SWIGINTERN
+#endif
+
+#ifndef SWIGRUNTIMEINLINE
+# define SWIGRUNTIMEINLINE SWIGRUNTIME SWIGINLINE
+#endif
+
+/*  Generic buffer size */
+#ifndef SWIG_BUFFER_SIZE
+# define SWIG_BUFFER_SIZE 1024
+#endif
+
+/* Flags for pointer conversions */
+#define SWIG_POINTER_DISOWN        0x1
+#define SWIG_CAST_NEW_MEMORY       0x2
+
+/* Flags for new pointer objects */
+#define SWIG_POINTER_OWN           0x1
+
+
+/* 
+   Flags/methods for returning states.
+   
+   The SWIG conversion methods, as ConvertPtr, return an integer 
+   that tells if the conversion was successful or not. And if not,
+   an error code can be returned (see swigerrors.swg for the codes).
+   
+   Use the following macros/flags to set or process the returning
+   states.
+   
+   In old versions of SWIG, code such as the following was usually written:
+
+     if (SWIG_ConvertPtr(obj,vptr,ty.flags) != -1) {
+       // success code
+     } else {
+       //fail code
+     }
+
+   Now you can be more explicit:
+
+    int res = SWIG_ConvertPtr(obj,vptr,ty.flags);
+    if (SWIG_IsOK(res)) {
+      // success code
+    } else {
+      // fail code
+    }
+
+   which is the same really, but now you can also do
+
+    Type *ptr;
+    int res = SWIG_ConvertPtr(obj,(void **)(&ptr),ty.flags);
+    if (SWIG_IsOK(res)) {
+      // success code
+      if (SWIG_IsNewObj(res) {
+        ...
+	delete *ptr;
+      } else {
+        ...
+      }
+    } else {
+      // fail code
+    }
+    
+   I.e., now SWIG_ConvertPtr can return new objects and you can
+   identify the case and take care of the deallocation. Of course that
+   also requires SWIG_ConvertPtr to return new result values, such as
+
+      int SWIG_ConvertPtr(obj, ptr,...) {         
+        if (<obj is ok>) {			       
+          if (<need new object>) {		       
+            *ptr = <ptr to new allocated object>; 
+            return SWIG_NEWOBJ;		       
+          } else {				       
+            *ptr = <ptr to old object>;	       
+            return SWIG_OLDOBJ;		       
+          } 				       
+        } else {				       
+          return SWIG_BADOBJ;		       
+        }					       
+      }
+
+   Of course, returning the plain '0(success)/-1(fail)' still works, but you can be
+   more explicit by returning SWIG_BADOBJ, SWIG_ERROR or any of the
+   SWIG errors code.
+
+   Finally, if the SWIG_CASTRANK_MODE is enabled, the result code
+   allows to return the 'cast rank', for example, if you have this
+
+       int food(double)
+       int fooi(int);
+
+   and you call
+ 
+      food(1)   // cast rank '1'  (1 -> 1.0)
+      fooi(1)   // cast rank '0'
+
+   just use the SWIG_AddCast()/SWIG_CheckState()
+*/
+
+#define SWIG_OK                    (0) 
+#define SWIG_ERROR                 (-1)
+#define SWIG_IsOK(r)               (r >= 0)
+#define SWIG_ArgError(r)           ((r != SWIG_ERROR) ? r : SWIG_TypeError)  
+
+/* The CastRankLimit says how many bits are used for the cast rank */
+#define SWIG_CASTRANKLIMIT         (1 << 8)
+/* The NewMask denotes the object was created (using new/malloc) */
+#define SWIG_NEWOBJMASK            (SWIG_CASTRANKLIMIT  << 1)
+/* The TmpMask is for in/out typemaps that use temporal objects */
+#define SWIG_TMPOBJMASK            (SWIG_NEWOBJMASK << 1)
+/* Simple returning values */
+#define SWIG_BADOBJ                (SWIG_ERROR)
+#define SWIG_OLDOBJ                (SWIG_OK)
+#define SWIG_NEWOBJ                (SWIG_OK | SWIG_NEWOBJMASK)
+#define SWIG_TMPOBJ                (SWIG_OK | SWIG_TMPOBJMASK)
+/* Check, add and del mask methods */
+#define SWIG_AddNewMask(r)         (SWIG_IsOK(r) ? (r | SWIG_NEWOBJMASK) : r)
+#define SWIG_DelNewMask(r)         (SWIG_IsOK(r) ? (r & ~SWIG_NEWOBJMASK) : r)
+#define SWIG_IsNewObj(r)           (SWIG_IsOK(r) && (r & SWIG_NEWOBJMASK))
+#define SWIG_AddTmpMask(r)         (SWIG_IsOK(r) ? (r | SWIG_TMPOBJMASK) : r)
+#define SWIG_DelTmpMask(r)         (SWIG_IsOK(r) ? (r & ~SWIG_TMPOBJMASK) : r)
+#define SWIG_IsTmpObj(r)           (SWIG_IsOK(r) && (r & SWIG_TMPOBJMASK))
+
+/* Cast-Rank Mode */
+#if defined(SWIG_CASTRANK_MODE)
+#  ifndef SWIG_TypeRank
+#    define SWIG_TypeRank             unsigned long
+#  endif
+#  ifndef SWIG_MAXCASTRANK            /* Default cast allowed */
+#    define SWIG_MAXCASTRANK          (2)
+#  endif
+#  define SWIG_CASTRANKMASK          ((SWIG_CASTRANKLIMIT) -1)
+#  define SWIG_CastRank(r)           (r & SWIG_CASTRANKMASK)
+SWIGINTERNINLINE int SWIG_AddCast(int r) { 
+  return SWIG_IsOK(r) ? ((SWIG_CastRank(r) < SWIG_MAXCASTRANK) ? (r + 1) : SWIG_ERROR) : r;
+}
+SWIGINTERNINLINE int SWIG_CheckState(int r) { 
+  return SWIG_IsOK(r) ? SWIG_CastRank(r) + 1 : 0; 
+}
+#else /* no cast-rank mode */
+#  define SWIG_AddCast
+#  define SWIG_CheckState(r) (SWIG_IsOK(r) ? 1 : 0)
+#endif
+
+
+#include <string.h>
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+typedef void *(*swig_converter_func)(void *, int *);
+typedef struct swig_type_info *(*swig_dycast_func)(void **);
+
+/* Structure to store information on one type */
+typedef struct swig_type_info {
+  const char             *name;			/* mangled name of this type */
+  const char             *str;			/* human readable name of this type */
+  swig_dycast_func        dcast;		/* dynamic cast function down a hierarchy */
+  struct swig_cast_info  *cast;			/* linked list of types that can cast into this type */
+  void                   *clientdata;		/* language specific type data */
+  int                    owndata;		/* flag if the structure owns the clientdata */
+} swig_type_info;
+
+/* Structure to store a type and conversion function used for casting */
+typedef struct swig_cast_info {
+  swig_type_info         *type;			/* pointer to type that is equivalent to this type */
+  swig_converter_func     converter;		/* function to cast the void pointers */
+  struct swig_cast_info  *next;			/* pointer to next cast in linked list */
+  struct swig_cast_info  *prev;			/* pointer to the previous cast */
+} swig_cast_info;
+
+/* Structure used to store module information
+ * Each module generates one structure like this, and the runtime collects
+ * all of these structures and stores them in a circularly linked list.*/
+typedef struct swig_module_info {
+  swig_type_info         **types;		/* Array of pointers to swig_type_info structures that are in this module */
+  size_t                 size;		        /* Number of types in this module */
+  struct swig_module_info *next;		/* Pointer to next element in circularly linked list */
+  swig_type_info         **type_initial;	/* Array of initially generated type structures */
+  swig_cast_info         **cast_initial;	/* Array of initially generated casting structures */
+  void                    *clientdata;		/* Language specific module data */
+} swig_module_info;
+
+/* 
+  Compare two type names skipping the space characters, therefore
+  "char*" == "char *" and "Class<int>" == "Class<int >", etc.
+
+  Return 0 when the two name types are equivalent, as in
+  strncmp, but skipping ' '.
+*/
+SWIGRUNTIME int
+SWIG_TypeNameComp(const char *f1, const char *l1,
+		  const char *f2, const char *l2) {
+  for (;(f1 != l1) && (f2 != l2); ++f1, ++f2) {
+    while ((*f1 == ' ') && (f1 != l1)) ++f1;
+    while ((*f2 == ' ') && (f2 != l2)) ++f2;
+    if (*f1 != *f2) return (*f1 > *f2) ? 1 : -1;
+  }
+  return (int)((l1 - f1) - (l2 - f2));
+}
+
+/*
+  Check type equivalence in a name list like <name1>|<name2>|...
+  Return 0 if not equal, 1 if equal
+*/
+SWIGRUNTIME int
+SWIG_TypeEquiv(const char *nb, const char *tb) {
+  int equiv = 0;
+  const char* te = tb + strlen(tb);
+  const char* ne = nb;
+  while (!equiv && *ne) {
+    for (nb = ne; *ne; ++ne) {
+      if (*ne == '|') break;
+    }
+    equiv = (SWIG_TypeNameComp(nb, ne, tb, te) == 0) ? 1 : 0;
+    if (*ne) ++ne;
+  }
+  return equiv;
+}
+
+/*
+  Check type equivalence in a name list like <name1>|<name2>|...
+  Return 0 if equal, -1 if nb < tb, 1 if nb > tb
+*/
+SWIGRUNTIME int
+SWIG_TypeCompare(const char *nb, const char *tb) {
+  int equiv = 0;
+  const char* te = tb + strlen(tb);
+  const char* ne = nb;
+  while (!equiv && *ne) {
+    for (nb = ne; *ne; ++ne) {
+      if (*ne == '|') break;
+    }
+    equiv = (SWIG_TypeNameComp(nb, ne, tb, te) == 0) ? 1 : 0;
+    if (*ne) ++ne;
+  }
+  return equiv;
+}
+
+
+/*
+  Check the typename
+*/
+SWIGRUNTIME swig_cast_info *
+SWIG_TypeCheck(const char *c, swig_type_info *ty) {
+  if (ty) {
+    swig_cast_info *iter = ty->cast;
+    while (iter) {
+      if (strcmp(iter->type->name, c) == 0) {
+        if (iter == ty->cast)
+          return iter;
+        /* Move iter to the top of the linked list */
+        iter->prev->next = iter->next;
+        if (iter->next)
+          iter->next->prev = iter->prev;
+        iter->next = ty->cast;
+        iter->prev = 0;
+        if (ty->cast) ty->cast->prev = iter;
+        ty->cast = iter;
+        return iter;
+      }
+      iter = iter->next;
+    }
+  }
+  return 0;
+}
+
+/* 
+  Identical to SWIG_TypeCheck, except strcmp is replaced with a pointer comparison
+*/
+SWIGRUNTIME swig_cast_info *
+SWIG_TypeCheckStruct(swig_type_info *from, swig_type_info *ty) {
+  if (ty) {
+    swig_cast_info *iter = ty->cast;
+    while (iter) {
+      if (iter->type == from) {
+        if (iter == ty->cast)
+          return iter;
+        /* Move iter to the top of the linked list */
+        iter->prev->next = iter->next;
+        if (iter->next)
+          iter->next->prev = iter->prev;
+        iter->next = ty->cast;
+        iter->prev = 0;
+        if (ty->cast) ty->cast->prev = iter;
+        ty->cast = iter;
+        return iter;
+      }
+      iter = iter->next;
+    }
+  }
+  return 0;
+}
+
+/*
+  Cast a pointer up an inheritance hierarchy
+*/
+SWIGRUNTIMEINLINE void *
+SWIG_TypeCast(swig_cast_info *ty, void *ptr, int *newmemory) {
+  return ((!ty) || (!ty->converter)) ? ptr : (*ty->converter)(ptr, newmemory);
+}
+
+/* 
+   Dynamic pointer casting. Down an inheritance hierarchy
+*/
+SWIGRUNTIME swig_type_info *
+SWIG_TypeDynamicCast(swig_type_info *ty, void **ptr) {
+  swig_type_info *lastty = ty;
+  if (!ty || !ty->dcast) return ty;
+  while (ty && (ty->dcast)) {
+    ty = (*ty->dcast)(ptr);
+    if (ty) lastty = ty;
+  }
+  return lastty;
+}
+
+/*
+  Return the name associated with this type
+*/
+SWIGRUNTIMEINLINE const char *
+SWIG_TypeName(const swig_type_info *ty) {
+  return ty->name;
+}
+
+/*
+  Return the pretty name associated with this type,
+  that is an unmangled type name in a form presentable to the user.
+*/
+SWIGRUNTIME const char *
+SWIG_TypePrettyName(const swig_type_info *type) {
+  /* The "str" field contains the equivalent pretty names of the
+     type, separated by vertical-bar characters.  We choose
+     to print the last name, as it is often (?) the most
+     specific. */
+  if (!type) return NULL;
+  if (type->str != NULL) {
+    const char *last_name = type->str;
+    const char *s;
+    for (s = type->str; *s; s++)
+      if (*s == '|') last_name = s+1;
+    return last_name;
+  }
+  else
+    return type->name;
+}
+
+/* 
+   Set the clientdata field for a type
+*/
+SWIGRUNTIME void
+SWIG_TypeClientData(swig_type_info *ti, void *clientdata) {
+  swig_cast_info *cast = ti->cast;
+  /* if (ti->clientdata == clientdata) return; */
+  ti->clientdata = clientdata;
+  
+  while (cast) {
+    if (!cast->converter) {
+      swig_type_info *tc = cast->type;
+      if (!tc->clientdata) {
+	SWIG_TypeClientData(tc, clientdata);
+      }
+    }    
+    cast = cast->next;
+  }
+}
+SWIGRUNTIME void
+SWIG_TypeNewClientData(swig_type_info *ti, void *clientdata) {
+  SWIG_TypeClientData(ti, clientdata);
+  ti->owndata = 1;
+}
+  
+/*
+  Search for a swig_type_info structure only by mangled name
+  Search is a O(log #types)
+  
+  We start searching at module start, and finish searching when start == end.  
+  Note: if start == end at the beginning of the function, we go all the way around
+  the circular list.
+*/
+SWIGRUNTIME swig_type_info *
+SWIG_MangledTypeQueryModule(swig_module_info *start, 
+                            swig_module_info *end, 
+		            const char *name) {
+  swig_module_info *iter = start;
+  do {
+    if (iter->size) {
+      register size_t l = 0;
+      register size_t r = iter->size - 1;
+      do {
+	/* since l+r >= 0, we can (>> 1) instead (/ 2) */
+	register size_t i = (l + r) >> 1; 
+	const char *iname = iter->types[i]->name;
+	if (iname) {
+	  register int compare = strcmp(name, iname);
+	  if (compare == 0) {	    
+	    return iter->types[i];
+	  } else if (compare < 0) {
+	    if (i) {
+	      r = i - 1;
+	    } else {
+	      break;
+	    }
+	  } else if (compare > 0) {
+	    l = i + 1;
+	  }
+	} else {
+	  break; /* should never happen */
+	}
+      } while (l <= r);
+    }
+    iter = iter->next;
+  } while (iter != end);
+  return 0;
+}
+
+/*
+  Search for a swig_type_info structure for either a mangled name or a human readable name.
+  It first searches the mangled names of the types, which is a O(log #types)
+  If a type is not found it then searches the human readable names, which is O(#types).
+  
+  We start searching at module start, and finish searching when start == end.  
+  Note: if start == end at the beginning of the function, we go all the way around
+  the circular list.
+*/
+SWIGRUNTIME swig_type_info *
+SWIG_TypeQueryModule(swig_module_info *start, 
+                     swig_module_info *end, 
+		     const char *name) {
+  /* STEP 1: Search the name field using binary search */
+  swig_type_info *ret = SWIG_MangledTypeQueryModule(start, end, name);
+  if (ret) {
+    return ret;
+  } else {
+    /* STEP 2: If the type hasn't been found, do a complete search
+       of the str field (the human readable name) */
+    swig_module_info *iter = start;
+    do {
+      register size_t i = 0;
+      for (; i < iter->size; ++i) {
+	if (iter->types[i]->str && (SWIG_TypeEquiv(iter->types[i]->str, name)))
+	  return iter->types[i];
+      }
+      iter = iter->next;
+    } while (iter != end);
+  }
+  
+  /* neither found a match */
+  return 0;
+}
+
+/* 
+   Pack binary data into a string
+*/
+SWIGRUNTIME char *
+SWIG_PackData(char *c, void *ptr, size_t sz) {
+  static const char hex[17] = "0123456789abcdef";
+  register const unsigned char *u = (unsigned char *) ptr;
+  register const unsigned char *eu =  u + sz;
+  for (; u != eu; ++u) {
+    register unsigned char uu = *u;
+    *(c++) = hex[(uu & 0xf0) >> 4];
+    *(c++) = hex[uu & 0xf];
+  }
+  return c;
+}
+
+/* 
+   Unpack binary data from a string
+*/
+SWIGRUNTIME const char *
+SWIG_UnpackData(const char *c, void *ptr, size_t sz) {
+  register unsigned char *u = (unsigned char *) ptr;
+  register const unsigned char *eu = u + sz;
+  for (; u != eu; ++u) {
+    register char d = *(c++);
+    register unsigned char uu;
+    if ((d >= '0') && (d <= '9'))
+      uu = ((d - '0') << 4);
+    else if ((d >= 'a') && (d <= 'f'))
+      uu = ((d - ('a'-10)) << 4);
+    else 
+      return (char *) 0;
+    d = *(c++);
+    if ((d >= '0') && (d <= '9'))
+      uu |= (d - '0');
+    else if ((d >= 'a') && (d <= 'f'))
+      uu |= (d - ('a'-10));
+    else 
+      return (char *) 0;
+    *u = uu;
+  }
+  return c;
+}
+
+/* 
+   Pack 'void *' into a string buffer.
+*/
+SWIGRUNTIME char *
+SWIG_PackVoidPtr(char *buff, void *ptr, const char *name, size_t bsz) {
+  char *r = buff;
+  if ((2*sizeof(void *) + 2) > bsz) return 0;
+  *(r++) = '_';
+  r = SWIG_PackData(r,&ptr,sizeof(void *));
+  if (strlen(name) + 1 > (bsz - (r - buff))) return 0;
+  strcpy(r,name);
+  return buff;
+}
+
+SWIGRUNTIME const char *
+SWIG_UnpackVoidPtr(const char *c, void **ptr, const char *name) {
+  if (*c != '_') {
+    if (strcmp(c,"NULL") == 0) {
+      *ptr = (void *) 0;
+      return name;
+    } else {
+      return 0;
+    }
+  }
+  return SWIG_UnpackData(++c,ptr,sizeof(void *));
+}
+
+SWIGRUNTIME char *
+SWIG_PackDataName(char *buff, void *ptr, size_t sz, const char *name, size_t bsz) {
+  char *r = buff;
+  size_t lname = (name ? strlen(name) : 0);
+  if ((2*sz + 2 + lname) > bsz) return 0;
+  *(r++) = '_';
+  r = SWIG_PackData(r,ptr,sz);
+  if (lname) {
+    strncpy(r,name,lname+1);
+  } else {
+    *r = 0;
+  }
+  return buff;
+}
+
+SWIGRUNTIME const char *
+SWIG_UnpackDataName(const char *c, void *ptr, size_t sz, const char *name) {
+  if (*c != '_') {
+    if (strcmp(c,"NULL") == 0) {
+      memset(ptr,0,sz);
+      return name;
+    } else {
+      return 0;
+    }
+  }
+  return SWIG_UnpackData(++c,ptr,sz);
+}
+
+#ifdef __cplusplus
+}
+#endif
+
+/*  Errors in SWIG */
+#define  SWIG_UnknownError    	   -1 
+#define  SWIG_IOError        	   -2 
+#define  SWIG_RuntimeError   	   -3 
+#define  SWIG_IndexError     	   -4 
+#define  SWIG_TypeError      	   -5 
+#define  SWIG_DivisionByZero 	   -6 
+#define  SWIG_OverflowError  	   -7 
+#define  SWIG_SyntaxError    	   -8 
+#define  SWIG_ValueError     	   -9 
+#define  SWIG_SystemError    	   -10
+#define  SWIG_AttributeError 	   -11
+#define  SWIG_MemoryError    	   -12 
+#define  SWIG_NullReferenceError   -13
+
+
+
+/* Compatibility macros for Python 3 */
+#if PY_VERSION_HEX >= 0x03000000
+
+#define PyClass_Check(obj) PyObject_IsInstance(obj, (PyObject *)&PyType_Type)
+#define PyInt_Check(x) PyLong_Check(x)
+#define PyInt_AsLong(x) PyLong_AsLong(x)
+#define PyInt_FromLong(x) PyLong_FromLong(x)
+#define PyInt_FromSize_t(x) PyLong_FromSize_t(x)
+#define PyString_Check(name) PyBytes_Check(name)
+#define PyString_FromString(x) PyUnicode_FromString(x)
+#define PyString_Format(fmt, args)  PyUnicode_Format(fmt, args)
+#define PyString_AsString(str) PyBytes_AsString(str)
+#define PyString_Size(str) PyBytes_Size(str)	
+#define PyString_InternFromString(key) PyUnicode_InternFromString(key)
+#define Py_TPFLAGS_HAVE_CLASS Py_TPFLAGS_BASETYPE
+#define PyString_AS_STRING(x) PyUnicode_AS_STRING(x)
+#define _PyLong_FromSsize_t(x) PyLong_FromSsize_t(x)
+
+#endif
+
+#ifndef Py_TYPE
+#  define Py_TYPE(op) ((op)->ob_type)
+#endif
+
+/* SWIG APIs for compatibility of both Python 2 & 3 */
+
+#if PY_VERSION_HEX >= 0x03000000
+#  define SWIG_Python_str_FromFormat PyUnicode_FromFormat
+#else
+#  define SWIG_Python_str_FromFormat PyString_FromFormat
+#endif
+
+
+/* Warning: This function will allocate a new string in Python 3,
+ * so please call SWIG_Python_str_DelForPy3(x) to free the space.
+ */
+SWIGINTERN char*
+SWIG_Python_str_AsChar(PyObject *str)
+{
+#if PY_VERSION_HEX >= 0x03000000
+  char *cstr;
+  char *newstr;
+  Py_ssize_t len;
+  str = PyUnicode_AsUTF8String(str);
+  PyBytes_AsStringAndSize(str, &cstr, &len);
+  newstr = (char *) malloc(len+1);
+  memcpy(newstr, cstr, len+1);
+  Py_XDECREF(str);
+  return newstr;
+#else
+  return PyString_AsString(str);
+#endif
+}
+
+#if PY_VERSION_HEX >= 0x03000000
+#  define SWIG_Python_str_DelForPy3(x) free( (void*) (x) )
+#else
+#  define SWIG_Python_str_DelForPy3(x) 
+#endif
+
+
+SWIGINTERN PyObject*
+SWIG_Python_str_FromChar(const char *c)
+{
+#if PY_VERSION_HEX >= 0x03000000
+  return PyUnicode_FromString(c); 
+#else
+  return PyString_FromString(c);
+#endif
+}
+
+/* Add PyOS_snprintf for old Pythons */
+#if PY_VERSION_HEX < 0x02020000
+# if defined(_MSC_VER) || defined(__BORLANDC__) || defined(_WATCOM)
+#  define PyOS_snprintf _snprintf
+# else
+#  define PyOS_snprintf snprintf
+# endif
+#endif
+
+/* A crude PyString_FromFormat implementation for old Pythons */
+#if PY_VERSION_HEX < 0x02020000
+
+#ifndef SWIG_PYBUFFER_SIZE
+# define SWIG_PYBUFFER_SIZE 1024
+#endif
+
+static PyObject *
+PyString_FromFormat(const char *fmt, ...) {
+  va_list ap;
+  char buf[SWIG_PYBUFFER_SIZE * 2];
+  int res;
+  va_start(ap, fmt);
+  res = vsnprintf(buf, sizeof(buf), fmt, ap);
+  va_end(ap);
+  return (res < 0 || res >= (int)sizeof(buf)) ? 0 : PyString_FromString(buf);
+}
+#endif
+
+/* Add PyObject_Del for old Pythons */
+#if PY_VERSION_HEX < 0x01060000
+# define PyObject_Del(op) PyMem_DEL((op))
+#endif
+#ifndef PyObject_DEL
+# define PyObject_DEL PyObject_Del
+#endif
+
+/* A crude PyExc_StopIteration exception for old Pythons */
+#if PY_VERSION_HEX < 0x02020000
+# ifndef PyExc_StopIteration
+#  define PyExc_StopIteration PyExc_RuntimeError
+# endif
+# ifndef PyObject_GenericGetAttr
+#  define PyObject_GenericGetAttr 0
+# endif
+#endif
+
+/* Py_NotImplemented is defined in 2.1 and up. */
+#if PY_VERSION_HEX < 0x02010000
+# ifndef Py_NotImplemented
+#  define Py_NotImplemented PyExc_RuntimeError
+# endif
+#endif
+
+/* A crude PyString_AsStringAndSize implementation for old Pythons */
+#if PY_VERSION_HEX < 0x02010000
+# ifndef PyString_AsStringAndSize
+#  define PyString_AsStringAndSize(obj, s, len) {*s = PyString_AsString(obj); *len = *s ? strlen(*s) : 0;}
+# endif
+#endif
+
+/* PySequence_Size for old Pythons */
+#if PY_VERSION_HEX < 0x02000000
+# ifndef PySequence_Size
+#  define PySequence_Size PySequence_Length
+# endif
+#endif
+
+/* PyBool_FromLong for old Pythons */
+#if PY_VERSION_HEX < 0x02030000
+static
+PyObject *PyBool_FromLong(long ok)
+{
+  PyObject *result = ok ? Py_True : Py_False;
+  Py_INCREF(result);
+  return result;
+}
+#endif
+
+/* Py_ssize_t for old Pythons */
+/* This code is as recommended by: */
+/* http://www.python.org/dev/peps/pep-0353/#conversion-guidelines */
+#if PY_VERSION_HEX < 0x02050000 && !defined(PY_SSIZE_T_MIN)
+typedef int Py_ssize_t;
+# define PY_SSIZE_T_MAX INT_MAX
+# define PY_SSIZE_T_MIN INT_MIN
+typedef inquiry lenfunc;
+typedef intargfunc ssizeargfunc;
+typedef intintargfunc ssizessizeargfunc;
+typedef intobjargproc ssizeobjargproc;
+typedef intintobjargproc ssizessizeobjargproc;
+typedef getreadbufferproc readbufferproc;
+typedef getwritebufferproc writebufferproc;
+typedef getsegcountproc segcountproc;
+typedef getcharbufferproc charbufferproc;
+static long PyNumber_AsSsize_t (PyObject *x, void *SWIGUNUSEDPARM(exc))
+{
+  long result = 0;
+  PyObject *i = PyNumber_Int(x);
+  if (i) {
+    result = PyInt_AsLong(i);
+    Py_DECREF(i);
+  }
+  return result;
+}
+#endif
+
+#if PY_VERSION_HEX < 0x02040000
+#define Py_VISIT(op)				\
+  do { 						\
+    if (op) {					\
+      int vret = visit((op), arg);		\
+      if (vret)					\
+        return vret;				\
+    }						\
+  } while (0)
+#endif
+
+#if PY_VERSION_HEX < 0x02030000
+typedef struct {
+  PyTypeObject type;
+  PyNumberMethods as_number;
+  PyMappingMethods as_mapping;
+  PySequenceMethods as_sequence;
+  PyBufferProcs as_buffer;
+  PyObject *name, *slots;
+} PyHeapTypeObject;
+#endif
+
+#if PY_VERSION_HEX < 0x02030000
+typedef destructor freefunc;
+#endif
+
+#if ((PY_MAJOR_VERSION == 2 && PY_MINOR_VERSION > 6) || \
+     (PY_MAJOR_VERSION == 3 && PY_MINOR_VERSION > 0) || \
+     (PY_MAJOR_VERSION > 3))
+# define SWIGPY_USE_CAPSULE
+# define SWIGPY_CAPSULE_NAME ((char*)"swig_runtime_data" SWIG_RUNTIME_VERSION ".type_pointer_capsule" SWIG_TYPE_TABLE_NAME)
+#endif
+
+#if PY_VERSION_HEX < 0x03020000
+#define PyDescr_TYPE(x) (((PyDescrObject *)(x))->d_type)
+#define PyDescr_NAME(x) (((PyDescrObject *)(x))->d_name)
+#endif
+
+/* -----------------------------------------------------------------------------
+ * error manipulation
+ * ----------------------------------------------------------------------------- */
+
+SWIGRUNTIME PyObject*
+SWIG_Python_ErrorType(int code) {
+  PyObject* type = 0;
+  switch(code) {
+  case SWIG_MemoryError:
+    type = PyExc_MemoryError;
+    break;
+  case SWIG_IOError:
+    type = PyExc_IOError;
+    break;
+  case SWIG_RuntimeError:
+    type = PyExc_RuntimeError;
+    break;
+  case SWIG_IndexError:
+    type = PyExc_IndexError;
+    break;
+  case SWIG_TypeError:
+    type = PyExc_TypeError;
+    break;
+  case SWIG_DivisionByZero:
+    type = PyExc_ZeroDivisionError;
+    break;
+  case SWIG_OverflowError:
+    type = PyExc_OverflowError;
+    break;
+  case SWIG_SyntaxError:
+    type = PyExc_SyntaxError;
+    break;
+  case SWIG_ValueError:
+    type = PyExc_ValueError;
+    break;
+  case SWIG_SystemError:
+    type = PyExc_SystemError;
+    break;
+  case SWIG_AttributeError:
+    type = PyExc_AttributeError;
+    break;
+  default:
+    type = PyExc_RuntimeError;
+  }
+  return type;
+}
+
+
+SWIGRUNTIME void
+SWIG_Python_AddErrorMsg(const char* mesg)
+{
+  PyObject *type = 0;
+  PyObject *value = 0;
+  PyObject *traceback = 0;
+
+  if (PyErr_Occurred()) PyErr_Fetch(&type, &value, &traceback);
+  if (value) {
+    char *tmp;
+    PyObject *old_str = PyObject_Str(value);
+    PyErr_Clear();
+    Py_XINCREF(type);
+
+    PyErr_Format(type, "%s %s", tmp = SWIG_Python_str_AsChar(old_str), mesg);
+    SWIG_Python_str_DelForPy3(tmp);
+    Py_DECREF(old_str);
+    Py_DECREF(value);
+  } else {
+    PyErr_SetString(PyExc_RuntimeError, mesg);
+  }
+}
+
+#if defined(SWIG_PYTHON_NO_THREADS)
+#  if defined(SWIG_PYTHON_THREADS)
+#    undef SWIG_PYTHON_THREADS
+#  endif
+#endif
+#if defined(SWIG_PYTHON_THREADS) /* Threading support is enabled */
+#  if !defined(SWIG_PYTHON_USE_GIL) && !defined(SWIG_PYTHON_NO_USE_GIL)
+#    if (PY_VERSION_HEX >= 0x02030000) /* For 2.3 or later, use the PyGILState calls */
+#      define SWIG_PYTHON_USE_GIL
+#    endif
+#  endif
+#  if defined(SWIG_PYTHON_USE_GIL) /* Use PyGILState threads calls */
+#    ifndef SWIG_PYTHON_INITIALIZE_THREADS
+#     define SWIG_PYTHON_INITIALIZE_THREADS  PyEval_InitThreads() 
+#    endif
+#    ifdef __cplusplus /* C++ code */
+       class SWIG_Python_Thread_Block {
+         bool status;
+         PyGILState_STATE state;
+       public:
+         void end() { if (status) { PyGILState_Release(state); status = false;} }
+         SWIG_Python_Thread_Block() : status(true), state(PyGILState_Ensure()) {}
+         ~SWIG_Python_Thread_Block() { end(); }
+       };
+       class SWIG_Python_Thread_Allow {
+         bool status;
+         PyThreadState *save;
+       public:
+         void end() { if (status) { PyEval_RestoreThread(save); status = false; }}
+         SWIG_Python_Thread_Allow() : status(true), save(PyEval_SaveThread()) {}
+         ~SWIG_Python_Thread_Allow() { end(); }
+       };
+#      define SWIG_PYTHON_THREAD_BEGIN_BLOCK   SWIG_Python_Thread_Block _swig_thread_block
+#      define SWIG_PYTHON_THREAD_END_BLOCK     _swig_thread_block.end()
+#      define SWIG_PYTHON_THREAD_BEGIN_ALLOW   SWIG_Python_Thread_Allow _swig_thread_allow
+#      define SWIG_PYTHON_THREAD_END_ALLOW     _swig_thread_allow.end()
+#    else /* C code */
+#      define SWIG_PYTHON_THREAD_BEGIN_BLOCK   PyGILState_STATE _swig_thread_block = PyGILState_Ensure()
+#      define SWIG_PYTHON_THREAD_END_BLOCK     PyGILState_Release(_swig_thread_block)
+#      define SWIG_PYTHON_THREAD_BEGIN_ALLOW   PyThreadState *_swig_thread_allow = PyEval_SaveThread()
+#      define SWIG_PYTHON_THREAD_END_ALLOW     PyEval_RestoreThread(_swig_thread_allow)
+#    endif
+#  else /* Old thread way, not implemented, user must provide it */
+#    if !defined(SWIG_PYTHON_INITIALIZE_THREADS)
+#      define SWIG_PYTHON_INITIALIZE_THREADS
+#    endif
+#    if !defined(SWIG_PYTHON_THREAD_BEGIN_BLOCK)
+#      define SWIG_PYTHON_THREAD_BEGIN_BLOCK
+#    endif
+#    if !defined(SWIG_PYTHON_THREAD_END_BLOCK)
+#      define SWIG_PYTHON_THREAD_END_BLOCK
+#    endif
+#    if !defined(SWIG_PYTHON_THREAD_BEGIN_ALLOW)
+#      define SWIG_PYTHON_THREAD_BEGIN_ALLOW
+#    endif
+#    if !defined(SWIG_PYTHON_THREAD_END_ALLOW)
+#      define SWIG_PYTHON_THREAD_END_ALLOW
+#    endif
+#  endif
+#else /* No thread support */
+#  define SWIG_PYTHON_INITIALIZE_THREADS
+#  define SWIG_PYTHON_THREAD_BEGIN_BLOCK
+#  define SWIG_PYTHON_THREAD_END_BLOCK
+#  define SWIG_PYTHON_THREAD_BEGIN_ALLOW
+#  define SWIG_PYTHON_THREAD_END_ALLOW
+#endif
+
+/* -----------------------------------------------------------------------------
+ * Python API portion that goes into the runtime
+ * ----------------------------------------------------------------------------- */
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/* -----------------------------------------------------------------------------
+ * Constant declarations
+ * ----------------------------------------------------------------------------- */
+
+/* Constant Types */
+#define SWIG_PY_POINTER 4
+#define SWIG_PY_BINARY  5
+
+/* Constant information structure */
+typedef struct swig_const_info {
+  int type;
+  char *name;
+  long lvalue;
+  double dvalue;
+  void   *pvalue;
+  swig_type_info **ptype;
+} swig_const_info;
+
+
+/* -----------------------------------------------------------------------------
+ * Wrapper of PyInstanceMethod_New() used in Python 3
+ * It is exported to the generated module, used for -fastproxy
+ * ----------------------------------------------------------------------------- */
+#if PY_VERSION_HEX >= 0x03000000
+SWIGRUNTIME PyObject* SWIG_PyInstanceMethod_New(PyObject *SWIGUNUSEDPARM(self), PyObject *func)
+{
+  return PyInstanceMethod_New(func);
+}
+#else
+SWIGRUNTIME PyObject* SWIG_PyInstanceMethod_New(PyObject *SWIGUNUSEDPARM(self), PyObject *SWIGUNUSEDPARM(func))
+{
+  return NULL;
+}
+#endif
+
+#ifdef __cplusplus
+}
+#endif
+
+
+/* -----------------------------------------------------------------------------
+ * pyrun.swg
+ *
+ * This file contains the runtime support for Python modules
+ * and includes code for managing global variables and pointer
+ * type checking.
+ *
+ * ----------------------------------------------------------------------------- */
+
+/* Common SWIG API */
+
+/* for raw pointers */
+#define SWIG_Python_ConvertPtr(obj, pptr, type, flags)  SWIG_Python_ConvertPtrAndOwn(obj, pptr, type, flags, 0)
+#define SWIG_ConvertPtr(obj, pptr, type, flags)         SWIG_Python_ConvertPtr(obj, pptr, type, flags)
+#define SWIG_ConvertPtrAndOwn(obj,pptr,type,flags,own)  SWIG_Python_ConvertPtrAndOwn(obj, pptr, type, flags, own)
+
+#ifdef SWIGPYTHON_BUILTIN
+#define SWIG_NewPointerObj(ptr, type, flags)            SWIG_Python_NewPointerObj(self, ptr, type, flags)
+#else
+#define SWIG_NewPointerObj(ptr, type, flags)            SWIG_Python_NewPointerObj(NULL, ptr, type, flags)
+#endif
+
+#define SWIG_InternalNewPointerObj(ptr, type, flags)	SWIG_Python_NewPointerObj(NULL, ptr, type, flags)
+
+#define SWIG_CheckImplicit(ty)                          SWIG_Python_CheckImplicit(ty) 
+#define SWIG_AcquirePtr(ptr, src)                       SWIG_Python_AcquirePtr(ptr, src)
+#define swig_owntype                                    int
+
+/* for raw packed data */
+#define SWIG_ConvertPacked(obj, ptr, sz, ty)            SWIG_Python_ConvertPacked(obj, ptr, sz, ty)
+#define SWIG_NewPackedObj(ptr, sz, type)                SWIG_Python_NewPackedObj(ptr, sz, type)
+
+/* for class or struct pointers */
+#define SWIG_ConvertInstance(obj, pptr, type, flags)    SWIG_ConvertPtr(obj, pptr, type, flags)
+#define SWIG_NewInstanceObj(ptr, type, flags)           SWIG_NewPointerObj(ptr, type, flags)
+
+/* for C or C++ function pointers */
+#define SWIG_ConvertFunctionPtr(obj, pptr, type)        SWIG_Python_ConvertFunctionPtr(obj, pptr, type)
+#define SWIG_NewFunctionPtrObj(ptr, type)               SWIG_Python_NewPointerObj(NULL, ptr, type, 0)
+
+/* for C++ member pointers, ie, member methods */
+#define SWIG_ConvertMember(obj, ptr, sz, ty)            SWIG_Python_ConvertPacked(obj, ptr, sz, ty)
+#define SWIG_NewMemberObj(ptr, sz, type)                SWIG_Python_NewPackedObj(ptr, sz, type)
+
+
+/* Runtime API */
+
+#define SWIG_GetModule(clientdata)                      SWIG_Python_GetModule()
+#define SWIG_SetModule(clientdata, pointer)             SWIG_Python_SetModule(pointer)
+#define SWIG_NewClientData(obj)                         SwigPyClientData_New(obj)
+
+#define SWIG_SetErrorObj                                SWIG_Python_SetErrorObj                            
+#define SWIG_SetErrorMsg                        	SWIG_Python_SetErrorMsg				   
+#define SWIG_ErrorType(code)                    	SWIG_Python_ErrorType(code)                        
+#define SWIG_Error(code, msg)            		SWIG_Python_SetErrorMsg(SWIG_ErrorType(code), msg) 
+#define SWIG_fail                        		goto fail					   
+
+
+/* Runtime API implementation */
+
+/* Error manipulation */
+
+SWIGINTERN void 
+SWIG_Python_SetErrorObj(PyObject *errtype, PyObject *obj) {
+  SWIG_PYTHON_THREAD_BEGIN_BLOCK; 
+  PyErr_SetObject(errtype, obj);
+  Py_DECREF(obj);
+  SWIG_PYTHON_THREAD_END_BLOCK;
+}
+
+SWIGINTERN void 
+SWIG_Python_SetErrorMsg(PyObject *errtype, const char *msg) {
+  SWIG_PYTHON_THREAD_BEGIN_BLOCK;
+  PyErr_SetString(errtype, msg);
+  SWIG_PYTHON_THREAD_END_BLOCK;
+}
+
+#define SWIG_Python_Raise(obj, type, desc)  SWIG_Python_SetErrorObj(SWIG_Python_ExceptionType(desc), obj)
+
+/* Set a constant value */
+
+#if defined(SWIGPYTHON_BUILTIN)
+
+SWIGINTERN void
+SwigPyBuiltin_AddPublicSymbol(PyObject *seq, const char *key) {
+  PyObject *s = PyString_InternFromString(key);
+  PyList_Append(seq, s);
+  Py_DECREF(s);
+}
+
+SWIGINTERN void
+SWIG_Python_SetConstant(PyObject *d, PyObject *public_interface, const char *name, PyObject *obj) {   
+#if PY_VERSION_HEX < 0x02030000
+  PyDict_SetItemString(d, (char *)name, obj);
+#else
+  PyDict_SetItemString(d, name, obj);
+#endif
+  Py_DECREF(obj);
+  if (public_interface)
+    SwigPyBuiltin_AddPublicSymbol(public_interface, name);
+}
+
+#else
+
+SWIGINTERN void
+SWIG_Python_SetConstant(PyObject *d, const char *name, PyObject *obj) {   
+#if PY_VERSION_HEX < 0x02030000
+  PyDict_SetItemString(d, (char *)name, obj);
+#else
+  PyDict_SetItemString(d, name, obj);
+#endif
+  Py_DECREF(obj);                            
+}
+
+#endif
+
+/* Append a value to the result obj */
+
+SWIGINTERN PyObject*
+SWIG_Python_AppendOutput(PyObject* result, PyObject* obj) {
+#if !defined(SWIG_PYTHON_OUTPUT_TUPLE)
+  if (!result) {
+    result = obj;
+  } else if (result == Py_None) {
+    Py_DECREF(result);
+    result = obj;
+  } else {
+    if (!PyList_Check(result)) {
+      PyObject *o2 = result;
+      result = PyList_New(1);
+      PyList_SetItem(result, 0, o2);
+    }
+    PyList_Append(result,obj);
+    Py_DECREF(obj);
+  }
+  return result;
+#else
+  PyObject*   o2;
+  PyObject*   o3;
+  if (!result) {
+    result = obj;
+  } else if (result == Py_None) {
+    Py_DECREF(result);
+    result = obj;
+  } else {
+    if (!PyTuple_Check(result)) {
+      o2 = result;
+      result = PyTuple_New(1);
+      PyTuple_SET_ITEM(result, 0, o2);
+    }
+    o3 = PyTuple_New(1);
+    PyTuple_SET_ITEM(o3, 0, obj);
+    o2 = result;
+    result = PySequence_Concat(o2, o3);
+    Py_DECREF(o2);
+    Py_DECREF(o3);
+  }
+  return result;
+#endif
+}
+
+/* Unpack the argument tuple */
+
+SWIGINTERN int
+SWIG_Python_UnpackTuple(PyObject *args, const char *name, Py_ssize_t min, Py_ssize_t max, PyObject **objs)
+{
+  if (!args) {
+    if (!min && !max) {
+      return 1;
+    } else {
+      PyErr_Format(PyExc_TypeError, "%s expected %s%d arguments, got none", 
+		   name, (min == max ? "" : "at least "), (int)min);
+      return 0;
+    }
+  }  
+  if (!PyTuple_Check(args)) {
+    if (min <= 1 && max >= 1) {
+      register int i;
+      objs[0] = args;
+      for (i = 1; i < max; ++i) {
+	objs[i] = 0;
+      }
+      return 2;
+    }
+    PyErr_SetString(PyExc_SystemError, "UnpackTuple() argument list is not a tuple");
+    return 0;
+  } else {
+    register Py_ssize_t l = PyTuple_GET_SIZE(args);
+    if (l < min) {
+      PyErr_Format(PyExc_TypeError, "%s expected %s%d arguments, got %d", 
+		   name, (min == max ? "" : "at least "), (int)min, (int)l);
+      return 0;
+    } else if (l > max) {
+      PyErr_Format(PyExc_TypeError, "%s expected %s%d arguments, got %d", 
+		   name, (min == max ? "" : "at most "), (int)max, (int)l);
+      return 0;
+    } else {
+      register int i;
+      for (i = 0; i < l; ++i) {
+	objs[i] = PyTuple_GET_ITEM(args, i);
+      }
+      for (; l < max; ++l) {
+	objs[l] = 0;
+      }
+      return i + 1;
+    }    
+  }
+}
+
+/* A functor is a function object with one single object argument */
+#if PY_VERSION_HEX >= 0x02020000
+#define SWIG_Python_CallFunctor(functor, obj)	        PyObject_CallFunctionObjArgs(functor, obj, NULL);
+#else
+#define SWIG_Python_CallFunctor(functor, obj)	        PyObject_CallFunction(functor, "O", obj);
+#endif
+
+/*
+  Helper for static pointer initialization for both C and C++ code, for example
+  static PyObject *SWIG_STATIC_POINTER(MyVar) = NewSomething(...);
+*/
+#ifdef __cplusplus
+#define SWIG_STATIC_POINTER(var)  var
+#else
+#define SWIG_STATIC_POINTER(var)  var = 0; if (!var) var
+#endif
+
+/* -----------------------------------------------------------------------------
+ * Pointer declarations
+ * ----------------------------------------------------------------------------- */
+
+/* Flags for new pointer objects */
+#define SWIG_POINTER_NOSHADOW       (SWIG_POINTER_OWN      << 1)
+#define SWIG_POINTER_NEW            (SWIG_POINTER_NOSHADOW | SWIG_POINTER_OWN)
+
+#define SWIG_POINTER_IMPLICIT_CONV  (SWIG_POINTER_DISOWN   << 1)
+
+#define SWIG_BUILTIN_TP_INIT	    (SWIG_POINTER_OWN << 2)
+#define SWIG_BUILTIN_INIT	    (SWIG_BUILTIN_TP_INIT | SWIG_POINTER_OWN)
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+/*  How to access Py_None */
+#if defined(_WIN32) || defined(__WIN32__) || defined(__CYGWIN__)
+#  ifndef SWIG_PYTHON_NO_BUILD_NONE
+#    ifndef SWIG_PYTHON_BUILD_NONE
+#      define SWIG_PYTHON_BUILD_NONE
+#    endif
+#  endif
+#endif
+
+#ifdef SWIG_PYTHON_BUILD_NONE
+#  ifdef Py_None
+#   undef Py_None
+#   define Py_None SWIG_Py_None()
+#  endif
+SWIGRUNTIMEINLINE PyObject * 
+_SWIG_Py_None(void)
+{
+  PyObject *none = Py_BuildValue((char*)"");
+  Py_DECREF(none);
+  return none;
+}
+SWIGRUNTIME PyObject * 
+SWIG_Py_None(void)
+{
+  static PyObject *SWIG_STATIC_POINTER(none) = _SWIG_Py_None();
+  return none;
+}
+#endif
+
+/* The python void return value */
+
+SWIGRUNTIMEINLINE PyObject * 
+SWIG_Py_Void(void)
+{
+  PyObject *none = Py_None;
+  Py_INCREF(none);
+  return none;
+}
+
+/* SwigPyClientData */
+
+typedef struct {
+  PyObject *klass;
+  PyObject *newraw;
+  PyObject *newargs;
+  PyObject *destroy;
+  int delargs;
+  int implicitconv;
+  PyTypeObject *pytype;
+} SwigPyClientData;
+
+SWIGRUNTIMEINLINE int 
+SWIG_Python_CheckImplicit(swig_type_info *ty)
+{
+  SwigPyClientData *data = (SwigPyClientData *)ty->clientdata;
+  return data ? data->implicitconv : 0;
+}
+
+SWIGRUNTIMEINLINE PyObject *
+SWIG_Python_ExceptionType(swig_type_info *desc) {
+  SwigPyClientData *data = desc ? (SwigPyClientData *) desc->clientdata : 0;
+  PyObject *klass = data ? data->klass : 0;
+  return (klass ? klass : PyExc_RuntimeError);
+}
+
+
+SWIGRUNTIME SwigPyClientData * 
+SwigPyClientData_New(PyObject* obj)
+{
+  if (!obj) {
+    return 0;
+  } else {
+    SwigPyClientData *data = (SwigPyClientData *)malloc(sizeof(SwigPyClientData));
+    /* the klass element */
+    data->klass = obj;
+    Py_INCREF(data->klass);
+    /* the newraw method and newargs arguments used to create a new raw instance */
+    if (PyClass_Check(obj)) {
+      data->newraw = 0;
+      data->newargs = obj;
+      Py_INCREF(obj);
+    } else {
+#if (PY_VERSION_HEX < 0x02020000)
+      data->newraw = 0;
+#else
+      data->newraw = PyObject_GetAttrString(data->klass, (char *)"__new__");
+#endif
+      if (data->newraw) {
+	Py_INCREF(data->newraw);
+	data->newargs = PyTuple_New(1);
+	PyTuple_SetItem(data->newargs, 0, obj);
+      } else {
+	data->newargs = obj;
+      }
+      Py_INCREF(data->newargs);
+    }
+    /* the destroy method, aka as the C++ delete method */
+    data->destroy = PyObject_GetAttrString(data->klass, (char *)"__swig_destroy__");
+    if (PyErr_Occurred()) {
+      PyErr_Clear();
+      data->destroy = 0;
+    }
+    if (data->destroy) {
+      int flags;
+      Py_INCREF(data->destroy);
+      flags = PyCFunction_GET_FLAGS(data->destroy);
+#ifdef METH_O
+      data->delargs = !(flags & (METH_O));
+#else
+      data->delargs = 0;
+#endif
+    } else {
+      data->delargs = 0;
+    }
+    data->implicitconv = 0;
+    data->pytype = 0;
+    return data;
+  }
+}
+
+SWIGRUNTIME void 
+SwigPyClientData_Del(SwigPyClientData *data) {
+  Py_XDECREF(data->newraw);
+  Py_XDECREF(data->newargs);
+  Py_XDECREF(data->destroy);
+}
+
+/* =============== SwigPyObject =====================*/
+
+typedef struct {
+  PyObject_HEAD
+  void *ptr;
+  swig_type_info *ty;
+  int own;
+  PyObject *next;
+#ifdef SWIGPYTHON_BUILTIN
+  PyObject *dict;
+#endif
+} SwigPyObject;
+
+SWIGRUNTIME PyObject *
+SwigPyObject_long(SwigPyObject *v)
+{
+  return PyLong_FromVoidPtr(v->ptr);
+}
+
+SWIGRUNTIME PyObject *
+SwigPyObject_format(const char* fmt, SwigPyObject *v)
+{
+  PyObject *res = NULL;
+  PyObject *args = PyTuple_New(1);
+  if (args) {
+    if (PyTuple_SetItem(args, 0, SwigPyObject_long(v)) == 0) {
+      PyObject *ofmt = SWIG_Python_str_FromChar(fmt);
+      if (ofmt) {
+#if PY_VERSION_HEX >= 0x03000000
+	res = PyUnicode_Format(ofmt,args);
+#else
+	res = PyString_Format(ofmt,args);
+#endif
+	Py_DECREF(ofmt);
+      }
+      Py_DECREF(args);
+    }
+  }
+  return res;
+}
+
+SWIGRUNTIME PyObject *
+SwigPyObject_oct(SwigPyObject *v)
+{
+  return SwigPyObject_format("%o",v);
+}
+
+SWIGRUNTIME PyObject *
+SwigPyObject_hex(SwigPyObject *v)
+{
+  return SwigPyObject_format("%x",v);
+}
+
+SWIGRUNTIME PyObject *
+#ifdef METH_NOARGS
+SwigPyObject_repr(SwigPyObject *v)
+#else
+SwigPyObject_repr(SwigPyObject *v, PyObject *args)
+#endif
+{
+  const char *name = SWIG_TypePrettyName(v->ty);
+  PyObject *repr = SWIG_Python_str_FromFormat("<Swig Object of type '%s' at %p>", (name ? name : "unknown"), (void *)v);
+  if (v->next) {
+# ifdef METH_NOARGS
+    PyObject *nrep = SwigPyObject_repr((SwigPyObject *)v->next);
+# else
+    PyObject *nrep = SwigPyObject_repr((SwigPyObject *)v->next, args);
+# endif
+# if PY_VERSION_HEX >= 0x03000000
+    PyObject *joined = PyUnicode_Concat(repr, nrep);
+    Py_DecRef(repr);
+    Py_DecRef(nrep);
+    repr = joined;
+# else
+    PyString_ConcatAndDel(&repr,nrep);
+# endif
+  }
+  return repr;  
+}
+
+SWIGRUNTIME int
+SwigPyObject_print(SwigPyObject *v, FILE *fp, int SWIGUNUSEDPARM(flags))
+{
+  char *str;
+#ifdef METH_NOARGS
+  PyObject *repr = SwigPyObject_repr(v);
+#else
+  PyObject *repr = SwigPyObject_repr(v, NULL);
+#endif
+  if (repr) {
+    str = SWIG_Python_str_AsChar(repr); 
+    fputs(str, fp);
+    SWIG_Python_str_DelForPy3(str);
+    Py_DECREF(repr);
+    return 0; 
+  } else {
+    return 1; 
+  }
+}
+
+SWIGRUNTIME PyObject *
+SwigPyObject_str(SwigPyObject *v)
+{
+  char result[SWIG_BUFFER_SIZE];
+  return SWIG_PackVoidPtr(result, v->ptr, v->ty->name, sizeof(result)) ?
+    SWIG_Python_str_FromChar(result) : 0;
+}
+
+SWIGRUNTIME int
+SwigPyObject_compare(SwigPyObject *v, SwigPyObject *w)
+{
+  void *i = v->ptr;
+  void *j = w->ptr;
+  return (i < j) ? -1 : ((i > j) ? 1 : 0);
+}
+
+/* Added for Python 3.x, would it also be useful for Python 2.x? */
+SWIGRUNTIME PyObject*
+SwigPyObject_richcompare(SwigPyObject *v, SwigPyObject *w, int op)
+{
+  PyObject* res;
+  if( op != Py_EQ && op != Py_NE ) {
+    Py_INCREF(Py_NotImplemented);
+    return Py_NotImplemented;
+  }
+  res = PyBool_FromLong( (SwigPyObject_compare(v, w)==0) == (op == Py_EQ) ? 1 : 0);
+  return res;  
+}
+
+
+SWIGRUNTIME PyTypeObject* SwigPyObject_TypeOnce(void);
+
+#ifdef SWIGPYTHON_BUILTIN
+static swig_type_info *SwigPyObject_stype = 0;
+SWIGRUNTIME PyTypeObject*
+SwigPyObject_type(void) {
+    SwigPyClientData *cd;
+    assert(SwigPyObject_stype);
+    cd = (SwigPyClientData*) SwigPyObject_stype->clientdata;
+    assert(cd);
+    assert(cd->pytype);
+    return cd->pytype;
+}
+#else
+SWIGRUNTIME PyTypeObject*
+SwigPyObject_type(void) {
+  static PyTypeObject *SWIG_STATIC_POINTER(type) = SwigPyObject_TypeOnce();
+  return type;
+}
+#endif
+
+SWIGRUNTIMEINLINE int
+SwigPyObject_Check(PyObject *op) {
+#ifdef SWIGPYTHON_BUILTIN
+  PyTypeObject *target_tp = SwigPyObject_type();
+  if (PyType_IsSubtype(op->ob_type, target_tp))
+    return 1;
+  return (strcmp(op->ob_type->tp_name, "SwigPyObject") == 0);
+#else
+  return (Py_TYPE(op) == SwigPyObject_type())
+    || (strcmp(Py_TYPE(op)->tp_name,"SwigPyObject") == 0);
+#endif
+}
+
+SWIGRUNTIME PyObject *
+SwigPyObject_New(void *ptr, swig_type_info *ty, int own);
+
+SWIGRUNTIME void
+SwigPyObject_dealloc(PyObject *v)
+{
+  SwigPyObject *sobj = (SwigPyObject *) v;
+  PyObject *next = sobj->next;
+  if (sobj->own == SWIG_POINTER_OWN) {
+    swig_type_info *ty = sobj->ty;
+    SwigPyClientData *data = ty ? (SwigPyClientData *) ty->clientdata : 0;
+    PyObject *destroy = data ? data->destroy : 0;
+    if (destroy) {
+      /* destroy is always a VARARGS method */
+      PyObject *res;
+      if (data->delargs) {
+	/* we need to create a temporary object to carry the destroy operation */
+	PyObject *tmp = SwigPyObject_New(sobj->ptr, ty, 0);
+	res = SWIG_Python_CallFunctor(destroy, tmp);
+	Py_DECREF(tmp);
+      } else {
+	PyCFunction meth = PyCFunction_GET_FUNCTION(destroy);
+	PyObject *mself = PyCFunction_GET_SELF(destroy);
+	res = ((*meth)(mself, v));
+      }
+      Py_XDECREF(res);
+    } 
+#if !defined(SWIG_PYTHON_SILENT_MEMLEAK)
+    else {
+      const char *name = SWIG_TypePrettyName(ty);
+      printf("swig/python detected a memory leak of type '%s', no destructor found.\n", (name ? name : "unknown"));
+    }
+#endif
+  } 
+  Py_XDECREF(next);
+  PyObject_DEL(v);
+}
+
+SWIGRUNTIME PyObject* 
+SwigPyObject_append(PyObject* v, PyObject* next)
+{
+  SwigPyObject *sobj = (SwigPyObject *) v;
+#ifndef METH_O
+  PyObject *tmp = 0;
+  if (!PyArg_ParseTuple(next,(char *)"O:append", &tmp)) return NULL;
+  next = tmp;
+#endif
+  if (!SwigPyObject_Check(next)) {
+    return NULL;
+  }
+  sobj->next = next;
+  Py_INCREF(next);
+  return SWIG_Py_Void();
+}
+
+SWIGRUNTIME PyObject* 
+#ifdef METH_NOARGS
+SwigPyObject_next(PyObject* v)
+#else
+SwigPyObject_next(PyObject* v, PyObject *SWIGUNUSEDPARM(args))
+#endif
+{
+  SwigPyObject *sobj = (SwigPyObject *) v;
+  if (sobj->next) {    
+    Py_INCREF(sobj->next);
+    return sobj->next;
+  } else {
+    return SWIG_Py_Void();
+  }
+}
+
+SWIGINTERN PyObject*
+#ifdef METH_NOARGS
+SwigPyObject_disown(PyObject *v)
+#else
+SwigPyObject_disown(PyObject* v, PyObject *SWIGUNUSEDPARM(args))
+#endif
+{
+  SwigPyObject *sobj = (SwigPyObject *)v;
+  sobj->own = 0;
+  return SWIG_Py_Void();
+}
+
+SWIGINTERN PyObject*
+#ifdef METH_NOARGS
+SwigPyObject_acquire(PyObject *v)
+#else
+SwigPyObject_acquire(PyObject* v, PyObject *SWIGUNUSEDPARM(args))
+#endif
+{
+  SwigPyObject *sobj = (SwigPyObject *)v;
+  sobj->own = SWIG_POINTER_OWN;
+  return SWIG_Py_Void();
+}
+
+SWIGINTERN PyObject*
+SwigPyObject_own(PyObject *v, PyObject *args)
+{
+  PyObject *val = 0;
+#if (PY_VERSION_HEX < 0x02020000)
+  if (!PyArg_ParseTuple(args,(char *)"|O:own",&val))
+#elif (PY_VERSION_HEX < 0x02050000)
+  if (!PyArg_UnpackTuple(args, (char *)"own", 0, 1, &val)) 
+#else
+  if (!PyArg_UnpackTuple(args, "own", 0, 1, &val)) 
+#endif
+    {
+      return NULL;
+    } 
+  else
+    {
+      SwigPyObject *sobj = (SwigPyObject *)v;
+      PyObject *obj = PyBool_FromLong(sobj->own);
+      if (val) {
+#ifdef METH_NOARGS
+	if (PyObject_IsTrue(val)) {
+	  SwigPyObject_acquire(v);
+	} else {
+	  SwigPyObject_disown(v);
+	}
+#else
+	if (PyObject_IsTrue(val)) {
+	  SwigPyObject_acquire(v,args);
+	} else {
+	  SwigPyObject_disown(v,args);
+	}
+#endif
+      } 
+      return obj;
+    }
+}
+
+#ifdef METH_O
+static PyMethodDef
+swigobject_methods[] = {
+  {(char *)"disown",  (PyCFunction)SwigPyObject_disown,  METH_NOARGS,  (char *)"releases ownership of the pointer"},
+  {(char *)"acquire", (PyCFunction)SwigPyObject_acquire, METH_NOARGS,  (char *)"aquires ownership of the pointer"},
+  {(char *)"own",     (PyCFunction)SwigPyObject_own,     METH_VARARGS, (char *)"returns/sets ownership of the pointer"},
+  {(char *)"append",  (PyCFunction)SwigPyObject_append,  METH_O,       (char *)"appends another 'this' object"},
+  {(char *)"next",    (PyCFunction)SwigPyObject_next,    METH_NOARGS,  (char *)"returns the next 'this' object"},
+  {(char *)"__repr__",(PyCFunction)SwigPyObject_repr,    METH_NOARGS,  (char *)"returns object representation"},
+  {0, 0, 0, 0}  
+};
+#else
+static PyMethodDef
+swigobject_methods[] = {
+  {(char *)"disown",  (PyCFunction)SwigPyObject_disown,  METH_VARARGS,  (char *)"releases ownership of the pointer"},
+  {(char *)"acquire", (PyCFunction)SwigPyObject_acquire, METH_VARARGS,  (char *)"aquires ownership of the pointer"},
+  {(char *)"own",     (PyCFunction)SwigPyObject_own,     METH_VARARGS,  (char *)"returns/sets ownership of the pointer"},
+  {(char *)"append",  (PyCFunction)SwigPyObject_append,  METH_VARARGS,  (char *)"appends another 'this' object"},
+  {(char *)"next",    (PyCFunction)SwigPyObject_next,    METH_VARARGS,  (char *)"returns the next 'this' object"},
+  {(char *)"__repr__",(PyCFunction)SwigPyObject_repr,   METH_VARARGS,  (char *)"returns object representation"},
+  {0, 0, 0, 0}  
+};
+#endif
+
+#if PY_VERSION_HEX < 0x02020000
+SWIGINTERN PyObject *
+SwigPyObject_getattr(SwigPyObject *sobj,char *name)
+{
+  return Py_FindMethod(swigobject_methods, (PyObject *)sobj, name);
+}
+#endif
+
+SWIGRUNTIME PyTypeObject*
+SwigPyObject_TypeOnce(void) {
+  static char swigobject_doc[] = "Swig object carries a C/C++ instance pointer";
+
+  static PyNumberMethods SwigPyObject_as_number = {
+    (binaryfunc)0, /*nb_add*/
+    (binaryfunc)0, /*nb_subtract*/
+    (binaryfunc)0, /*nb_multiply*/
+    /* nb_divide removed in Python 3 */
+#if PY_VERSION_HEX < 0x03000000
+    (binaryfunc)0, /*nb_divide*/
+#endif
+    (binaryfunc)0, /*nb_remainder*/
+    (binaryfunc)0, /*nb_divmod*/
+    (ternaryfunc)0,/*nb_power*/
+    (unaryfunc)0,  /*nb_negative*/
+    (unaryfunc)0,  /*nb_positive*/
+    (unaryfunc)0,  /*nb_absolute*/
+    (inquiry)0,    /*nb_nonzero*/
+    0,		   /*nb_invert*/
+    0,		   /*nb_lshift*/
+    0,		   /*nb_rshift*/
+    0,		   /*nb_and*/
+    0,		   /*nb_xor*/
+    0,		   /*nb_or*/
+#if PY_VERSION_HEX < 0x03000000
+    0,   /*nb_coerce*/
+#endif
+    (unaryfunc)SwigPyObject_long, /*nb_int*/
+#if PY_VERSION_HEX < 0x03000000
+    (unaryfunc)SwigPyObject_long, /*nb_long*/
+#else
+    0, /*nb_reserved*/
+#endif
+    (unaryfunc)0,                 /*nb_float*/
+#if PY_VERSION_HEX < 0x03000000
+    (unaryfunc)SwigPyObject_oct,  /*nb_oct*/
+    (unaryfunc)SwigPyObject_hex,  /*nb_hex*/
+#endif
+#if PY_VERSION_HEX >= 0x03000000 /* 3.0 */
+    0,0,0,0,0,0,0,0,0,0,0,0,0,0,0 /* nb_inplace_add -> nb_index, nb_inplace_divide removed */
+#elif PY_VERSION_HEX >= 0x02050000 /* 2.5.0 */
+    0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0 /* nb_inplace_add -> nb_index */
+#elif PY_VERSION_HEX >= 0x02020000 /* 2.2.0 */
+    0,0,0,0,0,0,0,0,0,0,0,0,0,0,0 /* nb_inplace_add -> nb_inplace_true_divide */
+#elif PY_VERSION_HEX >= 0x02000000 /* 2.0.0 */
+    0,0,0,0,0,0,0,0,0,0,0 /* nb_inplace_add -> nb_inplace_or */
+#endif
+  };
+
+  static PyTypeObject swigpyobject_type;
+  static int type_init = 0;
+  if (!type_init) {
+    const PyTypeObject tmp = {
+      /* PyObject header changed in Python 3 */
+#if PY_VERSION_HEX >= 0x03000000
+      PyVarObject_HEAD_INIT(NULL, 0)
+#else
+      PyObject_HEAD_INIT(NULL)
+      0,                                    /* ob_size */
+#endif
+      (char *)"SwigPyObject",               /* tp_name */
+      sizeof(SwigPyObject),                 /* tp_basicsize */
+      0,                                    /* tp_itemsize */
+      (destructor)SwigPyObject_dealloc,     /* tp_dealloc */
+      (printfunc)SwigPyObject_print,        /* tp_print */
+#if PY_VERSION_HEX < 0x02020000
+      (getattrfunc)SwigPyObject_getattr,    /* tp_getattr */
+#else
+      (getattrfunc)0,                       /* tp_getattr */
+#endif
+      (setattrfunc)0,                       /* tp_setattr */
+#if PY_VERSION_HEX >= 0x03000000
+    0, /* tp_reserved in 3.0.1, tp_compare in 3.0.0 but not used */
+#else
+      (cmpfunc)SwigPyObject_compare,        /* tp_compare */
+#endif
+      (reprfunc)SwigPyObject_repr,          /* tp_repr */
+      &SwigPyObject_as_number,              /* tp_as_number */
+      0,                                    /* tp_as_sequence */
+      0,                                    /* tp_as_mapping */
+      (hashfunc)0,                          /* tp_hash */
+      (ternaryfunc)0,                       /* tp_call */
+      (reprfunc)SwigPyObject_str,           /* tp_str */
+      PyObject_GenericGetAttr,              /* tp_getattro */
+      0,                                    /* tp_setattro */
+      0,                                    /* tp_as_buffer */
+      Py_TPFLAGS_DEFAULT,                   /* tp_flags */
+      swigobject_doc,                       /* tp_doc */
+      0,                                    /* tp_traverse */
+      0,                                    /* tp_clear */
+      (richcmpfunc)SwigPyObject_richcompare,/* tp_richcompare */
+      0,                                    /* tp_weaklistoffset */
+#if PY_VERSION_HEX >= 0x02020000
+      0,                                    /* tp_iter */
+      0,                                    /* tp_iternext */
+      swigobject_methods,                   /* tp_methods */
+      0,                                    /* tp_members */
+      0,                                    /* tp_getset */
+      0,                                    /* tp_base */
+      0,                                    /* tp_dict */
+      0,                                    /* tp_descr_get */
+      0,                                    /* tp_descr_set */
+      0,                                    /* tp_dictoffset */
+      0,                                    /* tp_init */
+      0,                                    /* tp_alloc */
+      0,                                    /* tp_new */
+      0,                                    /* tp_free */
+      0,                                    /* tp_is_gc */
+      0,                                    /* tp_bases */
+      0,                                    /* tp_mro */
+      0,                                    /* tp_cache */
+      0,                                    /* tp_subclasses */
+      0,                                    /* tp_weaklist */
+#endif
+#if PY_VERSION_HEX >= 0x02030000
+      0,                                    /* tp_del */
+#endif
+#if PY_VERSION_HEX >= 0x02060000
+      0,                                    /* tp_version */
+#endif
+#ifdef COUNT_ALLOCS
+      0,0,0,0                               /* tp_alloc -> tp_next */
+#endif
+    };
+    swigpyobject_type = tmp;
+    type_init = 1;
+#if PY_VERSION_HEX < 0x02020000
+    swigpyobject_type.ob_type = &PyType_Type;
+#else
+    if (PyType_Ready(&swigpyobject_type) < 0)
+      return NULL;
+#endif
+  }
+  return &swigpyobject_type;
+}
+
+SWIGRUNTIME PyObject *
+SwigPyObject_New(void *ptr, swig_type_info *ty, int own)
+{
+  SwigPyObject *sobj = PyObject_NEW(SwigPyObject, SwigPyObject_type());
+  if (sobj) {
+    sobj->ptr  = ptr;
+    sobj->ty   = ty;
+    sobj->own  = own;
+    sobj->next = 0;
+  }
+  return (PyObject *)sobj;
+}
+
+/* -----------------------------------------------------------------------------
+ * Implements a simple Swig Packed type, and use it instead of string
+ * ----------------------------------------------------------------------------- */
+
+typedef struct {
+  PyObject_HEAD
+  void *pack;
+  swig_type_info *ty;
+  size_t size;
+} SwigPyPacked;
+
+SWIGRUNTIME int
+SwigPyPacked_print(SwigPyPacked *v, FILE *fp, int SWIGUNUSEDPARM(flags))
+{
+  char result[SWIG_BUFFER_SIZE];
+  fputs("<Swig Packed ", fp); 
+  if (SWIG_PackDataName(result, v->pack, v->size, 0, sizeof(result))) {
+    fputs("at ", fp); 
+    fputs(result, fp); 
+  }
+  fputs(v->ty->name,fp); 
+  fputs(">", fp);
+  return 0; 
+}
+  
+SWIGRUNTIME PyObject *
+SwigPyPacked_repr(SwigPyPacked *v)
+{
+  char result[SWIG_BUFFER_SIZE];
+  if (SWIG_PackDataName(result, v->pack, v->size, 0, sizeof(result))) {
+    return SWIG_Python_str_FromFormat("<Swig Packed at %s%s>", result, v->ty->name);
+  } else {
+    return SWIG_Python_str_FromFormat("<Swig Packed %s>", v->ty->name);
+  }  
+}
+
+SWIGRUNTIME PyObject *
+SwigPyPacked_str(SwigPyPacked *v)
+{
+  char result[SWIG_BUFFER_SIZE];
+  if (SWIG_PackDataName(result, v->pack, v->size, 0, sizeof(result))){
+    return SWIG_Python_str_FromFormat("%s%s", result, v->ty->name);
+  } else {
+    return SWIG_Python_str_FromChar(v->ty->name);
+  }  
+}
+
+SWIGRUNTIME int
+SwigPyPacked_compare(SwigPyPacked *v, SwigPyPacked *w)
+{
+  size_t i = v->size;
+  size_t j = w->size;
+  int s = (i < j) ? -1 : ((i > j) ? 1 : 0);
+  return s ? s : strncmp((char *)v->pack, (char *)w->pack, 2*v->size);
+}
+
+SWIGRUNTIME PyTypeObject* SwigPyPacked_TypeOnce(void);
+
+SWIGRUNTIME PyTypeObject*
+SwigPyPacked_type(void) {
+  static PyTypeObject *SWIG_STATIC_POINTER(type) = SwigPyPacked_TypeOnce();
+  return type;
+}
+
+SWIGRUNTIMEINLINE int
+SwigPyPacked_Check(PyObject *op) {
+  return ((op)->ob_type == SwigPyPacked_TypeOnce()) 
+    || (strcmp((op)->ob_type->tp_name,"SwigPyPacked") == 0);
+}
+
+SWIGRUNTIME void
+SwigPyPacked_dealloc(PyObject *v)
+{
+  if (SwigPyPacked_Check(v)) {
+    SwigPyPacked *sobj = (SwigPyPacked *) v;
+    free(sobj->pack);
+  }
+  PyObject_DEL(v);
+}
+
+SWIGRUNTIME PyTypeObject*
+SwigPyPacked_TypeOnce(void) {
+  static char swigpacked_doc[] = "Swig object carries a C/C++ instance pointer";
+  static PyTypeObject swigpypacked_type;
+  static int type_init = 0;
+  if (!type_init) {
+    const PyTypeObject tmp = {
+      /* PyObject header changed in Python 3 */
+#if PY_VERSION_HEX>=0x03000000
+      PyVarObject_HEAD_INIT(NULL, 0)
+#else
+      PyObject_HEAD_INIT(NULL)
+      0,                                    /* ob_size */
+#endif
+      (char *)"SwigPyPacked",               /* tp_name */
+      sizeof(SwigPyPacked),                 /* tp_basicsize */
+      0,                                    /* tp_itemsize */
+      (destructor)SwigPyPacked_dealloc,     /* tp_dealloc */
+      (printfunc)SwigPyPacked_print,        /* tp_print */
+      (getattrfunc)0,                       /* tp_getattr */
+      (setattrfunc)0,                       /* tp_setattr */
+#if PY_VERSION_HEX>=0x03000000
+      0, /* tp_reserved in 3.0.1 */
+#else
+      (cmpfunc)SwigPyPacked_compare,        /* tp_compare */
+#endif
+      (reprfunc)SwigPyPacked_repr,          /* tp_repr */
+      0,                                    /* tp_as_number */
+      0,                                    /* tp_as_sequence */
+      0,                                    /* tp_as_mapping */
+      (hashfunc)0,                          /* tp_hash */
+      (ternaryfunc)0,                       /* tp_call */
+      (reprfunc)SwigPyPacked_str,           /* tp_str */
+      PyObject_GenericGetAttr,              /* tp_getattro */
+      0,                                    /* tp_setattro */
+      0,                                    /* tp_as_buffer */
+      Py_TPFLAGS_DEFAULT,                   /* tp_flags */
+      swigpacked_doc,                       /* tp_doc */
+      0,                                    /* tp_traverse */
+      0,                                    /* tp_clear */
+      0,                                    /* tp_richcompare */
+      0,                                    /* tp_weaklistoffset */
+#if PY_VERSION_HEX >= 0x02020000
+      0,                                    /* tp_iter */
+      0,                                    /* tp_iternext */
+      0,                                    /* tp_methods */
+      0,                                    /* tp_members */
+      0,                                    /* tp_getset */
+      0,                                    /* tp_base */
+      0,                                    /* tp_dict */
+      0,                                    /* tp_descr_get */
+      0,                                    /* tp_descr_set */
+      0,                                    /* tp_dictoffset */
+      0,                                    /* tp_init */
+      0,                                    /* tp_alloc */
+      0,                                    /* tp_new */
+      0,                                    /* tp_free */
+      0,                                    /* tp_is_gc */
+      0,                                    /* tp_bases */
+      0,                                    /* tp_mro */
+      0,                                    /* tp_cache */
+      0,                                    /* tp_subclasses */
+      0,                                    /* tp_weaklist */
+#endif
+#if PY_VERSION_HEX >= 0x02030000
+      0,                                    /* tp_del */
+#endif
+#if PY_VERSION_HEX >= 0x02060000
+      0,                                    /* tp_version */
+#endif
+#ifdef COUNT_ALLOCS
+      0,0,0,0                               /* tp_alloc -> tp_next */
+#endif
+    };
+    swigpypacked_type = tmp;
+    type_init = 1;
+#if PY_VERSION_HEX < 0x02020000
+    swigpypacked_type.ob_type = &PyType_Type;
+#else
+    if (PyType_Ready(&swigpypacked_type) < 0)
+      return NULL;
+#endif
+  }
+  return &swigpypacked_type;
+}
+
+SWIGRUNTIME PyObject *
+SwigPyPacked_New(void *ptr, size_t size, swig_type_info *ty)
+{
+  SwigPyPacked *sobj = PyObject_NEW(SwigPyPacked, SwigPyPacked_type());
+  if (sobj) {
+    void *pack = malloc(size);
+    if (pack) {
+      memcpy(pack, ptr, size);
+      sobj->pack = pack;
+      sobj->ty   = ty;
+      sobj->size = size;
+    } else {
+      PyObject_DEL((PyObject *) sobj);
+      sobj = 0;
+    }
+  }
+  return (PyObject *) sobj;
+}
+
+SWIGRUNTIME swig_type_info *
+SwigPyPacked_UnpackData(PyObject *obj, void *ptr, size_t size)
+{
+  if (SwigPyPacked_Check(obj)) {
+    SwigPyPacked *sobj = (SwigPyPacked *)obj;
+    if (sobj->size != size) return 0;
+    memcpy(ptr, sobj->pack, size);
+    return sobj->ty;
+  } else {
+    return 0;
+  }
+}
+
+/* -----------------------------------------------------------------------------
+ * pointers/data manipulation
+ * ----------------------------------------------------------------------------- */
+
+SWIGRUNTIMEINLINE PyObject *
+_SWIG_This(void)
+{
+    return SWIG_Python_str_FromChar("this");
+}
+
+static PyObject *swig_this = NULL;
+
+SWIGRUNTIME PyObject *
+SWIG_This(void)
+{
+  if (swig_this == NULL)
+    swig_this = _SWIG_This();
+  return swig_this;
+}
+
+/* #define SWIG_PYTHON_SLOW_GETSET_THIS */
+
+/* TODO: I don't know how to implement the fast getset in Python 3 right now */
+#if PY_VERSION_HEX>=0x03000000
+#define SWIG_PYTHON_SLOW_GETSET_THIS 
+#endif
+
+SWIGRUNTIME SwigPyObject *
+SWIG_Python_GetSwigThis(PyObject *pyobj) 
+{
+  PyObject *obj;
+
+  if (SwigPyObject_Check(pyobj))
+    return (SwigPyObject *) pyobj;
+
+#ifdef SWIGPYTHON_BUILTIN
+  (void)obj;
+# ifdef PyWeakref_CheckProxy
+  if (PyWeakref_CheckProxy(pyobj)) {
+    pyobj = PyWeakref_GET_OBJECT(pyobj);
+    if (pyobj && SwigPyObject_Check(pyobj))
+      return (SwigPyObject*) pyobj;
+  }
+# endif
+  return NULL;
+#else
+
+  obj = 0;
+
+#if (!defined(SWIG_PYTHON_SLOW_GETSET_THIS) && (PY_VERSION_HEX >= 0x02030000))
+  if (PyInstance_Check(pyobj)) {
+    obj = _PyInstance_Lookup(pyobj, SWIG_This());      
+  } else {
+    PyObject **dictptr = _PyObject_GetDictPtr(pyobj);
+    if (dictptr != NULL) {
+      PyObject *dict = *dictptr;
+      obj = dict ? PyDict_GetItem(dict, SWIG_This()) : 0;
+    } else {
+#ifdef PyWeakref_CheckProxy
+      if (PyWeakref_CheckProxy(pyobj)) {
+	PyObject *wobj = PyWeakref_GET_OBJECT(pyobj);
+	return wobj ? SWIG_Python_GetSwigThis(wobj) : 0;
+      }
+#endif
+      obj = PyObject_GetAttr(pyobj,SWIG_This());
+      if (obj) {
+	Py_DECREF(obj);
+      } else {
+	if (PyErr_Occurred()) PyErr_Clear();
+	return 0;
+      }
+    }
+  }
+#else
+  obj = PyObject_GetAttr(pyobj,SWIG_This());
+  if (obj) {
+    Py_DECREF(obj);
+  } else {
+    if (PyErr_Occurred()) PyErr_Clear();
+    return 0;
+  }
+#endif
+  if (obj && !SwigPyObject_Check(obj)) {
+    /* a PyObject is called 'this', try to get the 'real this'
+       SwigPyObject from it */ 
+    return SWIG_Python_GetSwigThis(obj);
+  }
+  return (SwigPyObject *)obj;
+#endif
+}
+
+/* Acquire a pointer value */
+
+SWIGRUNTIME int
+SWIG_Python_AcquirePtr(PyObject *obj, int own) {
+  if (own == SWIG_POINTER_OWN) {
+    SwigPyObject *sobj = SWIG_Python_GetSwigThis(obj);
+    if (sobj) {
+      int oldown = sobj->own;
+      sobj->own = own;
+      return oldown;
+    }
+  }
+  return 0;
+}
+
+/* Convert a pointer value */
+
+SWIGRUNTIME int
+SWIG_Python_ConvertPtrAndOwn(PyObject *obj, void **ptr, swig_type_info *ty, int flags, int *own) {
+  int res;
+  SwigPyObject *sobj;
+
+  if (!obj)
+    return SWIG_ERROR;
+  if (obj == Py_None) {
+    if (ptr)
+      *ptr = 0;
+    return SWIG_OK;
+  }
+
+  res = SWIG_ERROR;
+
+  sobj = SWIG_Python_GetSwigThis(obj);
+  if (own)
+    *own = 0;
+  while (sobj) {
+    void *vptr = sobj->ptr;
+    if (ty) {
+      swig_type_info *to = sobj->ty;
+      if (to == ty) {
+        /* no type cast needed */
+        if (ptr) *ptr = vptr;
+        break;
+      } else {
+        swig_cast_info *tc = SWIG_TypeCheck(to->name,ty);
+        if (!tc) {
+          sobj = (SwigPyObject *)sobj->next;
+        } else {
+          if (ptr) {
+            int newmemory = 0;
+            *ptr = SWIG_TypeCast(tc,vptr,&newmemory);
+            if (newmemory == SWIG_CAST_NEW_MEMORY) {
+              assert(own); /* badly formed typemap which will lead to a memory leak - it must set and use own to delete *ptr */
+              if (own)
+                *own = *own | SWIG_CAST_NEW_MEMORY;
+            }
+          }
+          break;
+        }
+      }
+    } else {
+      if (ptr) *ptr = vptr;
+      break;
+    }
+  }
+  if (sobj) {
+    if (own)
+      *own = *own | sobj->own;
+    if (flags & SWIG_POINTER_DISOWN) {
+      sobj->own = 0;
+    }
+    res = SWIG_OK;
+  } else {
+    if (flags & SWIG_POINTER_IMPLICIT_CONV) {
+      SwigPyClientData *data = ty ? (SwigPyClientData *) ty->clientdata : 0;
+      if (data && !data->implicitconv) {
+        PyObject *klass = data->klass;
+        if (klass) {
+          PyObject *impconv;
+          data->implicitconv = 1; /* avoid recursion and call 'explicit' constructors*/
+          impconv = SWIG_Python_CallFunctor(klass, obj);
+          data->implicitconv = 0;
+          if (PyErr_Occurred()) {
+            PyErr_Clear();
+            impconv = 0;
+          }
+          if (impconv) {
+            SwigPyObject *iobj = SWIG_Python_GetSwigThis(impconv);
+            if (iobj) {
+              void *vptr;
+              res = SWIG_Python_ConvertPtrAndOwn((PyObject*)iobj, &vptr, ty, 0, 0);
+              if (SWIG_IsOK(res)) {
+                if (ptr) {
+                  *ptr = vptr;
+                  /* transfer the ownership to 'ptr' */
+                  iobj->own = 0;
+                  res = SWIG_AddCast(res);
+                  res = SWIG_AddNewMask(res);
+                } else {
+                  res = SWIG_AddCast(res);		    
+                }
+              }
+            }
+            Py_DECREF(impconv);
+          }
+        }
+      }
+    }
+  }
+  return res;
+}
+
+/* Convert a function ptr value */
+
+SWIGRUNTIME int
+SWIG_Python_ConvertFunctionPtr(PyObject *obj, void **ptr, swig_type_info *ty) {
+  if (!PyCFunction_Check(obj)) {
+    return SWIG_ConvertPtr(obj, ptr, ty, 0);
+  } else {
+    void *vptr = 0;
+    
+    /* here we get the method pointer for callbacks */
+    const char *doc = (((PyCFunctionObject *)obj) -> m_ml -> ml_doc);
+    const char *desc = doc ? strstr(doc, "swig_ptr: ") : 0;
+    if (desc)
+      desc = ty ? SWIG_UnpackVoidPtr(desc + 10, &vptr, ty->name) : 0;
+    if (!desc) 
+      return SWIG_ERROR;
+    if (ty) {
+      swig_cast_info *tc = SWIG_TypeCheck(desc,ty);
+      if (tc) {
+        int newmemory = 0;
+        *ptr = SWIG_TypeCast(tc,vptr,&newmemory);
+        assert(!newmemory); /* newmemory handling not yet implemented */
+      } else {
+        return SWIG_ERROR;
+      }
+    } else {
+      *ptr = vptr;
+    }
+    return SWIG_OK;
+  }
+}
+
+/* Convert a packed value value */
+
+SWIGRUNTIME int
+SWIG_Python_ConvertPacked(PyObject *obj, void *ptr, size_t sz, swig_type_info *ty) {
+  swig_type_info *to = SwigPyPacked_UnpackData(obj, ptr, sz);
+  if (!to) return SWIG_ERROR;
+  if (ty) {
+    if (to != ty) {
+      /* check type cast? */
+      swig_cast_info *tc = SWIG_TypeCheck(to->name,ty);
+      if (!tc) return SWIG_ERROR;
+    }
+  }
+  return SWIG_OK;
+}  
+
+/* -----------------------------------------------------------------------------
+ * Create a new pointer object
+ * ----------------------------------------------------------------------------- */
+
+/*
+  Create a new instance object, without calling __init__, and set the
+  'this' attribute.
+*/
+
+SWIGRUNTIME PyObject* 
+SWIG_Python_NewShadowInstance(SwigPyClientData *data, PyObject *swig_this)
+{
+#if (PY_VERSION_HEX >= 0x02020000)
+  PyObject *inst = 0;
+  PyObject *newraw = data->newraw;
+  if (newraw) {
+    inst = PyObject_Call(newraw, data->newargs, NULL);
+    if (inst) {
+#if !defined(SWIG_PYTHON_SLOW_GETSET_THIS)
+      PyObject **dictptr = _PyObject_GetDictPtr(inst);
+      if (dictptr != NULL) {
+	PyObject *dict = *dictptr;
+	if (dict == NULL) {
+	  dict = PyDict_New();
+	  *dictptr = dict;
+	  PyDict_SetItem(dict, SWIG_This(), swig_this);
+	}
+      }
+#else
+      PyObject *key = SWIG_This();
+      PyObject_SetAttr(inst, key, swig_this);
+#endif
+    }
+  } else {
+#if PY_VERSION_HEX >= 0x03000000
+    inst = PyBaseObject_Type.tp_new((PyTypeObject*) data->newargs, Py_None, Py_None);
+    PyObject_SetAttr(inst, SWIG_This(), swig_this);
+    Py_TYPE(inst)->tp_flags &= ~Py_TPFLAGS_VALID_VERSION_TAG;
+#else
+    PyObject *dict = PyDict_New();
+    PyDict_SetItem(dict, SWIG_This(), swig_this);
+    inst = PyInstance_NewRaw(data->newargs, dict);
+    Py_DECREF(dict);
+#endif
+  }
+  return inst;
+#else
+#if (PY_VERSION_HEX >= 0x02010000)
+  PyObject *inst;
+  PyObject *dict = PyDict_New();
+  PyDict_SetItem(dict, SWIG_This(), swig_this);
+  inst = PyInstance_NewRaw(data->newargs, dict);
+  Py_DECREF(dict);
+  return (PyObject *) inst;
+#else
+  PyInstanceObject *inst = PyObject_NEW(PyInstanceObject, &PyInstance_Type);
+  if (inst == NULL) {
+    return NULL;
+  }
+  inst->in_class = (PyClassObject *)data->newargs;
+  Py_INCREF(inst->in_class);
+  inst->in_dict = PyDict_New();
+  if (inst->in_dict == NULL) {
+    Py_DECREF(inst);
+    return NULL;
+  }
+#ifdef Py_TPFLAGS_HAVE_WEAKREFS
+  inst->in_weakreflist = NULL;
+#endif
+#ifdef Py_TPFLAGS_GC
+  PyObject_GC_Init(inst);
+#endif
+  PyDict_SetItem(inst->in_dict, SWIG_This(), swig_this);
+  return (PyObject *) inst;
+#endif
+#endif
+}
+
+SWIGRUNTIME void
+SWIG_Python_SetSwigThis(PyObject *inst, PyObject *swig_this)
+{
+ PyObject *dict;
+#if (PY_VERSION_HEX >= 0x02020000) && !defined(SWIG_PYTHON_SLOW_GETSET_THIS)
+ PyObject **dictptr = _PyObject_GetDictPtr(inst);
+ if (dictptr != NULL) {
+   dict = *dictptr;
+   if (dict == NULL) {
+     dict = PyDict_New();
+     *dictptr = dict;
+   }
+   PyDict_SetItem(dict, SWIG_This(), swig_this);
+   return;
+ }
+#endif
+ dict = PyObject_GetAttrString(inst, (char*)"__dict__");
+ PyDict_SetItem(dict, SWIG_This(), swig_this);
+ Py_DECREF(dict);
+} 
+
+
+SWIGINTERN PyObject *
+SWIG_Python_InitShadowInstance(PyObject *args) {
+  PyObject *obj[2];
+  if (!SWIG_Python_UnpackTuple(args, "swiginit", 2, 2, obj)) {
+    return NULL;
+  } else {
+    SwigPyObject *sthis = SWIG_Python_GetSwigThis(obj[0]);
+    if (sthis) {
+      SwigPyObject_append((PyObject*) sthis, obj[1]);
+    } else {
+      SWIG_Python_SetSwigThis(obj[0], obj[1]);
+    }
+    return SWIG_Py_Void();
+  }
+}
+
+/* Create a new pointer object */
+
+SWIGRUNTIME PyObject *
+SWIG_Python_NewPointerObj(PyObject *self, void *ptr, swig_type_info *type, int flags) {
+  SwigPyClientData *clientdata;
+  PyObject * robj;
+  int own;
+
+  if (!ptr)
+    return SWIG_Py_Void();
+
+  clientdata = type ? (SwigPyClientData *)(type->clientdata) : 0;
+  own = (flags & SWIG_POINTER_OWN) ? SWIG_POINTER_OWN : 0;
+  if (clientdata && clientdata->pytype) {
+    SwigPyObject *newobj;
+    if (flags & SWIG_BUILTIN_TP_INIT) {
+      newobj = (SwigPyObject*) self;
+      if (newobj->ptr) {
+        PyObject *next_self = clientdata->pytype->tp_alloc(clientdata->pytype, 0);
+        while (newobj->next)
+	  newobj = (SwigPyObject *) newobj->next;
+        newobj->next = next_self;
+        newobj = (SwigPyObject *)next_self;
+      }
+    } else {
+      newobj = PyObject_New(SwigPyObject, clientdata->pytype);
+    }
+    if (newobj) {
+      newobj->ptr = ptr;
+      newobj->ty = type;
+      newobj->own = own;
+      newobj->next = 0;
+#ifdef SWIGPYTHON_BUILTIN
+      newobj->dict = 0;
+#endif
+      return (PyObject*) newobj;
+    }
+    return SWIG_Py_Void();
+  }
+
+  assert(!(flags & SWIG_BUILTIN_TP_INIT));
+
+  robj = SwigPyObject_New(ptr, type, own);
+  if (clientdata && !(flags & SWIG_POINTER_NOSHADOW)) {
+    PyObject *inst = SWIG_Python_NewShadowInstance(clientdata, robj);
+    if (inst) {
+      Py_DECREF(robj);
+      robj = inst;
+    }
+  }
+  return robj;
+}
+
+/* Create a new packed object */
+
+SWIGRUNTIMEINLINE PyObject *
+SWIG_Python_NewPackedObj(void *ptr, size_t sz, swig_type_info *type) {
+  return ptr ? SwigPyPacked_New((void *) ptr, sz, type) : SWIG_Py_Void();
+}
+
+/* -----------------------------------------------------------------------------*
+ *  Get type list 
+ * -----------------------------------------------------------------------------*/
+
+#ifdef SWIG_LINK_RUNTIME
+void *SWIG_ReturnGlobalTypeList(void *);
+#endif
+
+SWIGRUNTIME swig_module_info *
+SWIG_Python_GetModule(void) {
+  static void *type_pointer = (void *)0;
+  /* first check if module already created */
+  if (!type_pointer) {
+#ifdef SWIG_LINK_RUNTIME
+    type_pointer = SWIG_ReturnGlobalTypeList((void *)0);
+#else
+# ifdef SWIGPY_USE_CAPSULE
+    type_pointer = PyCapsule_Import(SWIGPY_CAPSULE_NAME, 0);
+# else
+    type_pointer = PyCObject_Import((char*)"swig_runtime_data" SWIG_RUNTIME_VERSION,
+				    (char*)"type_pointer" SWIG_TYPE_TABLE_NAME);
+# endif
+    if (PyErr_Occurred()) {
+      PyErr_Clear();
+      type_pointer = (void *)0;
+    }
+#endif
+  }
+  return (swig_module_info *) type_pointer;
+}
+
+#if PY_MAJOR_VERSION < 2
+/* PyModule_AddObject function was introduced in Python 2.0.  The following function
+   is copied out of Python/modsupport.c in python version 2.3.4 */
+SWIGINTERN int
+PyModule_AddObject(PyObject *m, char *name, PyObject *o)
+{
+  PyObject *dict;
+  if (!PyModule_Check(m)) {
+    PyErr_SetString(PyExc_TypeError,
+		    "PyModule_AddObject() needs module as first arg");
+    return SWIG_ERROR;
+  }
+  if (!o) {
+    PyErr_SetString(PyExc_TypeError,
+		    "PyModule_AddObject() needs non-NULL value");
+    return SWIG_ERROR;
+  }
+  
+  dict = PyModule_GetDict(m);
+  if (dict == NULL) {
+    /* Internal error -- modules must have a dict! */
+    PyErr_Format(PyExc_SystemError, "module '%s' has no __dict__",
+		 PyModule_GetName(m));
+    return SWIG_ERROR;
+  }
+  if (PyDict_SetItemString(dict, name, o))
+    return SWIG_ERROR;
+  Py_DECREF(o);
+  return SWIG_OK;
+}
+#endif
+
+SWIGRUNTIME void
+#ifdef SWIGPY_USE_CAPSULE
+SWIG_Python_DestroyModule(PyObject *obj)
+#else
+SWIG_Python_DestroyModule(void *vptr)
+#endif
+{
+#ifdef SWIGPY_USE_CAPSULE
+  swig_module_info *swig_module = (swig_module_info *) PyCapsule_GetPointer(obj, SWIGPY_CAPSULE_NAME);
+#else
+  swig_module_info *swig_module = (swig_module_info *) vptr;
+#endif
+  swig_type_info **types = swig_module->types;
+  size_t i;
+  for (i =0; i < swig_module->size; ++i) {
+    swig_type_info *ty = types[i];
+    if (ty->owndata) {
+      SwigPyClientData *data = (SwigPyClientData *) ty->clientdata;
+      if (data) SwigPyClientData_Del(data);
+    }
+  }
+  Py_DECREF(SWIG_This());
+  swig_this = NULL;
+}
+
+SWIGRUNTIME void
+SWIG_Python_SetModule(swig_module_info *swig_module) {
+#if PY_VERSION_HEX >= 0x03000000
+ /* Add a dummy module object into sys.modules */
+  PyObject *module = PyImport_AddModule((char*)"swig_runtime_data" SWIG_RUNTIME_VERSION);
+#else
+  static PyMethodDef swig_empty_runtime_method_table[] = { {NULL, NULL, 0, NULL} }; /* Sentinel */
+  PyObject *module = Py_InitModule((char*)"swig_runtime_data" SWIG_RUNTIME_VERSION, swig_empty_runtime_method_table);
+#endif
+#ifdef SWIGPY_USE_CAPSULE
+  PyObject *pointer = PyCapsule_New((void *) swig_module, SWIGPY_CAPSULE_NAME, SWIG_Python_DestroyModule);
+  if (pointer && module) {
+    PyModule_AddObject(module, (char*)"type_pointer_capsule" SWIG_TYPE_TABLE_NAME, pointer);
+  } else {
+    Py_XDECREF(pointer);
+  }
+#else
+  PyObject *pointer = PyCObject_FromVoidPtr((void *) swig_module, SWIG_Python_DestroyModule);
+  if (pointer && module) {
+    PyModule_AddObject(module, (char*)"type_pointer" SWIG_TYPE_TABLE_NAME, pointer);
+  } else {
+    Py_XDECREF(pointer);
+  }
+#endif
+}
+
+/* The python cached type query */
+SWIGRUNTIME PyObject *
+SWIG_Python_TypeCache(void) {
+  static PyObject *SWIG_STATIC_POINTER(cache) = PyDict_New();
+  return cache;
+}
+
+SWIGRUNTIME swig_type_info *
+SWIG_Python_TypeQuery(const char *type)
+{
+  PyObject *cache = SWIG_Python_TypeCache();
+  PyObject *key = SWIG_Python_str_FromChar(type); 
+  PyObject *obj = PyDict_GetItem(cache, key);
+  swig_type_info *descriptor;
+  if (obj) {
+#ifdef SWIGPY_USE_CAPSULE
+    descriptor = (swig_type_info *) PyCapsule_GetPointer(obj, NULL);
+#else
+    descriptor = (swig_type_info *) PyCObject_AsVoidPtr(obj);
+#endif
+  } else {
+    swig_module_info *swig_module = SWIG_Python_GetModule();
+    descriptor = SWIG_TypeQueryModule(swig_module, swig_module, type);
+    if (descriptor) {
+#ifdef SWIGPY_USE_CAPSULE
+      obj = PyCapsule_New((void*) descriptor, NULL, NULL);
+#else
+      obj = PyCObject_FromVoidPtr(descriptor, NULL);
+#endif
+      PyDict_SetItem(cache, key, obj);
+      Py_DECREF(obj);
+    }
+  }
+  Py_DECREF(key);
+  return descriptor;
+}
+
+/* 
+   For backward compatibility only
+*/
+#define SWIG_POINTER_EXCEPTION  0
+#define SWIG_arg_fail(arg)      SWIG_Python_ArgFail(arg)
+#define SWIG_MustGetPtr(p, type, argnum, flags)  SWIG_Python_MustGetPtr(p, type, argnum, flags)
+
+SWIGRUNTIME int
+SWIG_Python_AddErrMesg(const char* mesg, int infront)
+{  
+  if (PyErr_Occurred()) {
+    PyObject *type = 0;
+    PyObject *value = 0;
+    PyObject *traceback = 0;
+    PyErr_Fetch(&type, &value, &traceback);
+    if (value) {
+      char *tmp;
+      PyObject *old_str = PyObject_Str(value);
+      Py_XINCREF(type);
+      PyErr_Clear();
+      if (infront) {
+	PyErr_Format(type, "%s %s", mesg, tmp = SWIG_Python_str_AsChar(old_str));
+      } else {
+	PyErr_Format(type, "%s %s", tmp = SWIG_Python_str_AsChar(old_str), mesg);
+      }
+      SWIG_Python_str_DelForPy3(tmp);
+      Py_DECREF(old_str);
+    }
+    return 1;
+  } else {
+    return 0;
+  }
+}
+  
+SWIGRUNTIME int
+SWIG_Python_ArgFail(int argnum)
+{
+  if (PyErr_Occurred()) {
+    /* add information about failing argument */
+    char mesg[256];
+    PyOS_snprintf(mesg, sizeof(mesg), "argument number %d:", argnum);
+    return SWIG_Python_AddErrMesg(mesg, 1);
+  } else {
+    return 0;
+  }
+}
+
+SWIGRUNTIMEINLINE const char *
+SwigPyObject_GetDesc(PyObject *self)
+{
+  SwigPyObject *v = (SwigPyObject *)self;
+  swig_type_info *ty = v ? v->ty : 0;
+  return ty ? ty->str : "";
+}
+
+SWIGRUNTIME void
+SWIG_Python_TypeError(const char *type, PyObject *obj)
+{
+  if (type) {
+#if defined(SWIG_COBJECT_TYPES)
+    if (obj && SwigPyObject_Check(obj)) {
+      const char *otype = (const char *) SwigPyObject_GetDesc(obj);
+      if (otype) {
+	PyErr_Format(PyExc_TypeError, "a '%s' is expected, 'SwigPyObject(%s)' is received",
+		     type, otype);
+	return;
+      }
+    } else 
+#endif      
+    {
+      const char *otype = (obj ? obj->ob_type->tp_name : 0); 
+      if (otype) {
+	PyObject *str = PyObject_Str(obj);
+	const char *cstr = str ? SWIG_Python_str_AsChar(str) : 0;
+	if (cstr) {
+	  PyErr_Format(PyExc_TypeError, "a '%s' is expected, '%s(%s)' is received",
+		       type, otype, cstr);
+          SWIG_Python_str_DelForPy3(cstr);
+	} else {
+	  PyErr_Format(PyExc_TypeError, "a '%s' is expected, '%s' is received",
+		       type, otype);
+	}
+	Py_XDECREF(str);
+	return;
+      }
+    }   
+    PyErr_Format(PyExc_TypeError, "a '%s' is expected", type);
+  } else {
+    PyErr_Format(PyExc_TypeError, "unexpected type is received");
+  }
+}
+
+
+/* Convert a pointer value, signal an exception on a type mismatch */
+SWIGRUNTIME void *
+SWIG_Python_MustGetPtr(PyObject *obj, swig_type_info *ty, int SWIGUNUSEDPARM(argnum), int flags) {
+  void *result;
+  if (SWIG_Python_ConvertPtr(obj, &result, ty, flags) == -1) {
+    PyErr_Clear();
+#if SWIG_POINTER_EXCEPTION
+    if (flags) {
+      SWIG_Python_TypeError(SWIG_TypePrettyName(ty), obj);
+      SWIG_Python_ArgFail(argnum);
+    }
+#endif
+  }
+  return result;
+}
+
+#ifdef SWIGPYTHON_BUILTIN
+SWIGRUNTIME int
+SWIG_Python_NonDynamicSetAttr(PyObject *obj, PyObject *name, PyObject *value) {
+  PyTypeObject *tp = obj->ob_type;
+  PyObject *descr;
+  PyObject *encoded_name;
+  descrsetfunc f;
+  int res;
+
+# ifdef Py_USING_UNICODE
+  if (PyString_Check(name)) {
+    name = PyUnicode_Decode(PyString_AsString(name), PyString_Size(name), NULL, NULL);
+    if (!name)
+      return -1;
+  } else if (!PyUnicode_Check(name))
+# else
+  if (!PyString_Check(name))
+# endif
+  {
+    PyErr_Format(PyExc_TypeError, "attribute name must be string, not '%.200s'", name->ob_type->tp_name);
+    return -1;
+  } else {
+    Py_INCREF(name);
+  }
+
+  if (!tp->tp_dict) {
+    if (PyType_Ready(tp) < 0)
+      goto done;
+  }
+
+  res = -1;
+  descr = _PyType_Lookup(tp, name);
+  f = NULL;
+  if (descr != NULL)
+    f = descr->ob_type->tp_descr_set;
+  if (!f) {
+    if (PyString_Check(name)) {
+      encoded_name = name;
+      Py_INCREF(name);
+    } else {
+      encoded_name = PyUnicode_AsUTF8String(name);
+    }
+    PyErr_Format(PyExc_AttributeError, "'%.100s' object has no attribute '%.200s'", tp->tp_name, PyString_AsString(encoded_name));
+    Py_DECREF(encoded_name);
+  } else {
+    res = f(descr, obj, value);
+  }
+  
+  done:
+  Py_DECREF(name);
+  return res;
+}
+#endif
+
+
+#ifdef __cplusplus
+}
+#endif
+
+
+
+#define SWIG_exception_fail(code, msg) do { SWIG_Error(code, msg); SWIG_fail; } while(0) 
+
+#define SWIG_contract_assert(expr, msg) if (!(expr)) { SWIG_Error(SWIG_RuntimeError, msg); SWIG_fail; } else 
+
+
+
+/* -------- TYPES TABLE (BEGIN) -------- */
+
+#define SWIGTYPE_p_char swig_types[0]
+static swig_type_info *swig_types[2];
+static swig_module_info swig_module = {swig_types, 1, 0, 0, 0, 0};
+#define SWIG_TypeQuery(name) SWIG_TypeQueryModule(&swig_module, &swig_module, name)
+#define SWIG_MangledTypeQuery(name) SWIG_MangledTypeQueryModule(&swig_module, &swig_module, name)
+
+/* -------- TYPES TABLE (END) -------- */
+
+#if (PY_VERSION_HEX <= 0x02000000)
+# if !defined(SWIG_PYTHON_CLASSIC)
+#  error "This python version requires swig to be run with the '-classic' option"
+# endif
+#endif
+
+/*-----------------------------------------------
+              @(target):= _hello.so
+  ------------------------------------------------*/
+#if PY_VERSION_HEX >= 0x03000000
+#  define SWIG_init    PyInit__hello
+
+#else
+#  define SWIG_init    init_hello
+
+#endif
+#define SWIG_name    "_hello"
+
+#define SWIGVERSION 0x020007 
+#define SWIG_VERSION SWIGVERSION
+
+
+#define SWIG_as_voidptr(a) (void *)((const void *)(a)) 
+#define SWIG_as_voidptrptr(a) ((void)SWIG_as_voidptr(*a),(void**)(a)) 
+
+
+SWIGINTERN swig_type_info*
+SWIG_pchar_descriptor(void)
+{
+  static int init = 0;
+  static swig_type_info* info = 0;
+  if (!init) {
+    info = SWIG_TypeQuery("_p_char");
+    init = 1;
+  }
+  return info;
+}
+
+
+SWIGINTERN int
+SWIG_AsCharPtrAndSize(PyObject *obj, char** cptr, size_t* psize, int *alloc)
+{
+#if PY_VERSION_HEX>=0x03000000
+  if (PyUnicode_Check(obj))
+#else  
+  if (PyString_Check(obj))
+#endif
+  {
+    char *cstr; Py_ssize_t len;
+#if PY_VERSION_HEX>=0x03000000
+    if (!alloc && cptr) {
+        /* We can't allow converting without allocation, since the internal
+           representation of string in Python 3 is UCS-2/UCS-4 but we require
+           a UTF-8 representation.
+           TODO(bhy) More detailed explanation */
+        return SWIG_RuntimeError;
+    }
+    obj = PyUnicode_AsUTF8String(obj);
+    PyBytes_AsStringAndSize(obj, &cstr, &len);
+    if(alloc) *alloc = SWIG_NEWOBJ;
+#else
+    PyString_AsStringAndSize(obj, &cstr, &len);
+#endif
+    if (cptr) {
+      if (alloc) {
+	/* 
+	   In python the user should not be able to modify the inner
+	   string representation. To warranty that, if you define
+	   SWIG_PYTHON_SAFE_CSTRINGS, a new/copy of the python string
+	   buffer is always returned.
+
+	   The default behavior is just to return the pointer value,
+	   so, be careful.
+	*/ 
+#if defined(SWIG_PYTHON_SAFE_CSTRINGS)
+	if (*alloc != SWIG_OLDOBJ) 
+#else
+	if (*alloc == SWIG_NEWOBJ) 
+#endif
+	  {
+	    *cptr = (char *)memcpy((char *)malloc((len + 1)*sizeof(char)), cstr, sizeof(char)*(len + 1));
+	    *alloc = SWIG_NEWOBJ;
+	  }
+	else {
+	  *cptr = cstr;
+	  *alloc = SWIG_OLDOBJ;
+	}
+      } else {
+        #if PY_VERSION_HEX>=0x03000000
+        assert(0); /* Should never reach here in Python 3 */
+        #endif
+	*cptr = SWIG_Python_str_AsChar(obj);
+      }
+    }
+    if (psize) *psize = len + 1;
+#if PY_VERSION_HEX>=0x03000000
+    Py_XDECREF(obj);
+#endif
+    return SWIG_OK;
+  } else {
+    swig_type_info* pchar_descriptor = SWIG_pchar_descriptor();
+    if (pchar_descriptor) {
+      void* vptr = 0;
+      if (SWIG_ConvertPtr(obj, &vptr, pchar_descriptor, 0) == SWIG_OK) {
+	if (cptr) *cptr = (char *) vptr;
+	if (psize) *psize = vptr ? (strlen((char *)vptr) + 1) : 0;
+	if (alloc) *alloc = SWIG_OLDOBJ;
+	return SWIG_OK;
+      }
+    }
+  }
+  return SWIG_TypeError;
+}
+
+
+
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+SWIGINTERN PyObject *_wrap_say_hello(PyObject *SWIGUNUSEDPARM(self), PyObject *args) {
+  PyObject *resultobj = 0;
+  char *arg1 = (char *) 0 ;
+  int res1 ;
+  char *buf1 = 0 ;
+  int alloc1 = 0 ;
+  PyObject * obj0 = 0 ;
+  
+  if (!PyArg_ParseTuple(args,(char *)"O:say_hello",&obj0)) SWIG_fail;
+  res1 = SWIG_AsCharPtrAndSize(obj0, &buf1, NULL, &alloc1);
+  if (!SWIG_IsOK(res1)) {
+    SWIG_exception_fail(SWIG_ArgError(res1), "in method '" "say_hello" "', argument " "1"" of type '" "char const *""'");
+  }
+  arg1 = (char *)(buf1);
+  say_hello((char const *)arg1);
+  resultobj = SWIG_Py_Void();
+  if (alloc1 == SWIG_NEWOBJ) free((char*)buf1);
+  return resultobj;
+fail:
+  if (alloc1 == SWIG_NEWOBJ) free((char*)buf1);
+  return NULL;
+}
+
+
+static PyMethodDef SwigMethods[] = {
+	 { (char *)"SWIG_PyInstanceMethod_New", (PyCFunction)SWIG_PyInstanceMethod_New, METH_O, NULL},
+	 { (char *)"say_hello", _wrap_say_hello, METH_VARARGS, NULL},
+	 { NULL, NULL, 0, NULL }
+};
+
+
+/* -------- TYPE CONVERSION AND EQUIVALENCE RULES (BEGIN) -------- */
+
+static swig_type_info _swigt__p_char = {"_p_char", "char *", 0, 0, (void*)0, 0};
+
+static swig_type_info *swig_type_initial[] = {
+  &_swigt__p_char,
+};
+
+static swig_cast_info _swigc__p_char[] = {  {&_swigt__p_char, 0, 0, 0},{0, 0, 0, 0}};
+
+static swig_cast_info *swig_cast_initial[] = {
+  _swigc__p_char,
+};
+
+
+/* -------- TYPE CONVERSION AND EQUIVALENCE RULES (END) -------- */
+
+static swig_const_info swig_const_table[] = {
+{0, 0, 0, 0.0, 0, 0}};
+
+#ifdef __cplusplus
+}
+#endif
+/* -----------------------------------------------------------------------------
+ * Type initialization:
+ * This problem is tough by the requirement that no dynamic 
+ * memory is used. Also, since swig_type_info structures store pointers to 
+ * swig_cast_info structures and swig_cast_info structures store pointers back
+ * to swig_type_info structures, we need some lookup code at initialization. 
+ * The idea is that swig generates all the structures that are needed. 
+ * The runtime then collects these partially filled structures. 
+ * The SWIG_InitializeModule function takes these initial arrays out of 
+ * swig_module, and does all the lookup, filling in the swig_module.types
+ * array with the correct data and linking the correct swig_cast_info
+ * structures together.
+ *
+ * The generated swig_type_info structures are assigned staticly to an initial 
+ * array. We just loop through that array, and handle each type individually.
+ * First we lookup if this type has been already loaded, and if so, use the
+ * loaded structure instead of the generated one. Then we have to fill in the
+ * cast linked list. The cast data is initially stored in something like a
+ * two-dimensional array. Each row corresponds to a type (there are the same
+ * number of rows as there are in the swig_type_initial array). Each entry in
+ * a column is one of the swig_cast_info structures for that type.
+ * The cast_initial array is actually an array of arrays, because each row has
+ * a variable number of columns. So to actually build the cast linked list,
+ * we find the array of casts associated with the type, and loop through it 
+ * adding the casts to the list. The one last trick we need to do is making
+ * sure the type pointer in the swig_cast_info struct is correct.
+ *
+ * First off, we lookup the cast->type name to see if it is already loaded. 
+ * There are three cases to handle:
+ *  1) If the cast->type has already been loaded AND the type we are adding
+ *     casting info to has not been loaded (it is in this module), THEN we
+ *     replace the cast->type pointer with the type pointer that has already
+ *     been loaded.
+ *  2) If BOTH types (the one we are adding casting info to, and the 
+ *     cast->type) are loaded, THEN the cast info has already been loaded by
+ *     the previous module so we just ignore it.
+ *  3) Finally, if cast->type has not already been loaded, then we add that
+ *     swig_cast_info to the linked list (because the cast->type) pointer will
+ *     be correct.
+ * ----------------------------------------------------------------------------- */
+
+#ifdef __cplusplus
+extern "C" {
+#if 0
+} /* c-mode */
+#endif
+#endif
+
+#if 0
+#define SWIGRUNTIME_DEBUG
+#endif
+
+
+SWIGRUNTIME void
+SWIG_InitializeModule(void *clientdata) {
+  size_t i;
+  swig_module_info *module_head, *iter;
+  int found, init;
+  
+  clientdata = clientdata;
+  
+  /* check to see if the circular list has been setup, if not, set it up */
+  if (swig_module.next==0) {
+    /* Initialize the swig_module */
+    swig_module.type_initial = swig_type_initial;
+    swig_module.cast_initial = swig_cast_initial;
+    swig_module.next = &swig_module;
+    init = 1;
+  } else {
+    init = 0;
+  }
+  
+  /* Try and load any already created modules */
+  module_head = SWIG_GetModule(clientdata);
+  if (!module_head) {
+    /* This is the first module loaded for this interpreter */
+    /* so set the swig module into the interpreter */
+    SWIG_SetModule(clientdata, &swig_module);
+    module_head = &swig_module;
+  } else {
+    /* the interpreter has loaded a SWIG module, but has it loaded this one? */
+    found=0;
+    iter=module_head;
+    do {
+      if (iter==&swig_module) {
+        found=1;
+        break;
+      }
+      iter=iter->next;
+    } while (iter!= module_head);
+    
+    /* if the is found in the list, then all is done and we may leave */
+    if (found) return;
+    /* otherwise we must add out module into the list */
+    swig_module.next = module_head->next;
+    module_head->next = &swig_module;
+  }
+  
+  /* When multiple interpeters are used, a module could have already been initialized in
+       a different interpreter, but not yet have a pointer in this interpreter.
+       In this case, we do not want to continue adding types... everything should be
+       set up already */
+  if (init == 0) return;
+  
+  /* Now work on filling in swig_module.types */
+#ifdef SWIGRUNTIME_DEBUG
+  printf("SWIG_InitializeModule: size %d\n", swig_module.size);
+#endif
+  for (i = 0; i < swig_module.size; ++i) {
+    swig_type_info *type = 0;
+    swig_type_info *ret;
+    swig_cast_info *cast;
+    
+#ifdef SWIGRUNTIME_DEBUG
+    printf("SWIG_InitializeModule: type %d %s\n", i, swig_module.type_initial[i]->name);
+#endif
+    
+    /* if there is another module already loaded */
+    if (swig_module.next != &swig_module) {
+      type = SWIG_MangledTypeQueryModule(swig_module.next, &swig_module, swig_module.type_initial[i]->name);
+    }
+    if (type) {
+      /* Overwrite clientdata field */
+#ifdef SWIGRUNTIME_DEBUG
+      printf("SWIG_InitializeModule: found type %s\n", type->name);
+#endif
+      if (swig_module.type_initial[i]->clientdata) {
+        type->clientdata = swig_module.type_initial[i]->clientdata;
+#ifdef SWIGRUNTIME_DEBUG
+        printf("SWIG_InitializeModule: found and overwrite type %s \n", type->name);
+#endif
+      }
+    } else {
+      type = swig_module.type_initial[i];
+    }
+    
+    /* Insert casting types */
+    cast = swig_module.cast_initial[i];
+    while (cast->type) {
+      /* Don't need to add information already in the list */
+      ret = 0;
+#ifdef SWIGRUNTIME_DEBUG
+      printf("SWIG_InitializeModule: look cast %s\n", cast->type->name);
+#endif
+      if (swig_module.next != &swig_module) {
+        ret = SWIG_MangledTypeQueryModule(swig_module.next, &swig_module, cast->type->name);
+#ifdef SWIGRUNTIME_DEBUG
+        if (ret) printf("SWIG_InitializeModule: found cast %s\n", ret->name);
+#endif
+      }
+      if (ret) {
+        if (type == swig_module.type_initial[i]) {
+#ifdef SWIGRUNTIME_DEBUG
+          printf("SWIG_InitializeModule: skip old type %s\n", ret->name);
+#endif
+          cast->type = ret;
+          ret = 0;
+        } else {
+          /* Check for casting already in the list */
+          swig_cast_info *ocast = SWIG_TypeCheck(ret->name, type);
+#ifdef SWIGRUNTIME_DEBUG
+          if (ocast) printf("SWIG_InitializeModule: skip old cast %s\n", ret->name);
+#endif
+          if (!ocast) ret = 0;
+        }
+      }
+      
+      if (!ret) {
+#ifdef SWIGRUNTIME_DEBUG
+        printf("SWIG_InitializeModule: adding cast %s\n", cast->type->name);
+#endif
+        if (type->cast) {
+          type->cast->prev = cast;
+          cast->next = type->cast;
+        }
+        type->cast = cast;
+      }
+      cast++;
+    }
+    /* Set entry in modules->types array equal to the type */
+    swig_module.types[i] = type;
+  }
+  swig_module.types[i] = 0;
+  
+#ifdef SWIGRUNTIME_DEBUG
+  printf("**** SWIG_InitializeModule: Cast List ******\n");
+  for (i = 0; i < swig_module.size; ++i) {
+    int j = 0;
+    swig_cast_info *cast = swig_module.cast_initial[i];
+    printf("SWIG_InitializeModule: type %d %s\n", i, swig_module.type_initial[i]->name);
+    while (cast->type) {
+      printf("SWIG_InitializeModule: cast type %s\n", cast->type->name);
+      cast++;
+      ++j;
+    }
+    printf("---- Total casts: %d\n",j);
+  }
+  printf("**** SWIG_InitializeModule: Cast List ******\n");
+#endif
+}
+
+/* This function will propagate the clientdata field of type to
+* any new swig_type_info structures that have been added into the list
+* of equivalent types.  It is like calling
+* SWIG_TypeClientData(type, clientdata) a second time.
+*/
+SWIGRUNTIME void
+SWIG_PropagateClientData(void) {
+  size_t i;
+  swig_cast_info *equiv;
+  static int init_run = 0;
+  
+  if (init_run) return;
+  init_run = 1;
+  
+  for (i = 0; i < swig_module.size; i++) {
+    if (swig_module.types[i]->clientdata) {
+      equiv = swig_module.types[i]->cast;
+      while (equiv) {
+        if (!equiv->converter) {
+          if (equiv->type && !equiv->type->clientdata)
+          SWIG_TypeClientData(equiv->type, swig_module.types[i]->clientdata);
+        }
+        equiv = equiv->next;
+      }
+    }
+  }
+}
+
+#ifdef __cplusplus
+#if 0
+{
+  /* c-mode */
+#endif
+}
+#endif
+
+
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+  
+  /* Python-specific SWIG API */
+#define SWIG_newvarlink()                             SWIG_Python_newvarlink()
+#define SWIG_addvarlink(p, name, get_attr, set_attr)  SWIG_Python_addvarlink(p, name, get_attr, set_attr)
+#define SWIG_InstallConstants(d, constants)           SWIG_Python_InstallConstants(d, constants)
+  
+  /* -----------------------------------------------------------------------------
+   * global variable support code.
+   * ----------------------------------------------------------------------------- */
+  
+  typedef struct swig_globalvar {
+    char       *name;                  /* Name of global variable */
+    PyObject *(*get_attr)(void);       /* Return the current value */
+    int       (*set_attr)(PyObject *); /* Set the value */
+    struct swig_globalvar *next;
+  } swig_globalvar;
+  
+  typedef struct swig_varlinkobject {
+    PyObject_HEAD
+    swig_globalvar *vars;
+  } swig_varlinkobject;
+  
+  SWIGINTERN PyObject *
+  swig_varlink_repr(swig_varlinkobject *SWIGUNUSEDPARM(v)) {
+#if PY_VERSION_HEX >= 0x03000000
+    return PyUnicode_InternFromString("<Swig global variables>");
+#else
+    return PyString_FromString("<Swig global variables>");
+#endif
+  }
+  
+  SWIGINTERN PyObject *
+  swig_varlink_str(swig_varlinkobject *v) {
+#if PY_VERSION_HEX >= 0x03000000
+    PyObject *str = PyUnicode_InternFromString("(");
+    PyObject *tail;
+    PyObject *joined;
+    swig_globalvar *var;
+    for (var = v->vars; var; var=var->next) {
+      tail = PyUnicode_FromString(var->name);
+      joined = PyUnicode_Concat(str, tail);
+      Py_DecRef(str);
+      Py_DecRef(tail);
+      str = joined;
+      if (var->next) {
+        tail = PyUnicode_InternFromString(", ");
+        joined = PyUnicode_Concat(str, tail);
+        Py_DecRef(str);
+        Py_DecRef(tail);
+        str = joined;
+      }
+    }
+    tail = PyUnicode_InternFromString(")");
+    joined = PyUnicode_Concat(str, tail);
+    Py_DecRef(str);
+    Py_DecRef(tail);
+    str = joined;
+#else
+    PyObject *str = PyString_FromString("(");
+    swig_globalvar *var;
+    for (var = v->vars; var; var=var->next) {
+      PyString_ConcatAndDel(&str,PyString_FromString(var->name));
+      if (var->next) PyString_ConcatAndDel(&str,PyString_FromString(", "));
+    }
+    PyString_ConcatAndDel(&str,PyString_FromString(")"));
+#endif
+    return str;
+  }
+  
+  SWIGINTERN int
+  swig_varlink_print(swig_varlinkobject *v, FILE *fp, int SWIGUNUSEDPARM(flags)) {
+    char *tmp;
+    PyObject *str = swig_varlink_str(v);
+    fprintf(fp,"Swig global variables ");
+    fprintf(fp,"%s\n", tmp = SWIG_Python_str_AsChar(str));
+    SWIG_Python_str_DelForPy3(tmp);
+    Py_DECREF(str);
+    return 0;
+  }
+  
+  SWIGINTERN void
+  swig_varlink_dealloc(swig_varlinkobject *v) {
+    swig_globalvar *var = v->vars;
+    while (var) {
+      swig_globalvar *n = var->next;
+      free(var->name);
+      free(var);
+      var = n;
+    }
+  }
+  
+  SWIGINTERN PyObject *
+  swig_varlink_getattr(swig_varlinkobject *v, char *n) {
+    PyObject *res = NULL;
+    swig_globalvar *var = v->vars;
+    while (var) {
+      if (strcmp(var->name,n) == 0) {
+        res = (*var->get_attr)();
+        break;
+      }
+      var = var->next;
+    }
+    if (res == NULL && !PyErr_Occurred()) {
+      PyErr_SetString(PyExc_NameError,"Unknown C global variable");
+    }
+    return res;
+  }
+  
+  SWIGINTERN int
+  swig_varlink_setattr(swig_varlinkobject *v, char *n, PyObject *p) {
+    int res = 1;
+    swig_globalvar *var = v->vars;
+    while (var) {
+      if (strcmp(var->name,n) == 0) {
+        res = (*var->set_attr)(p);
+        break;
+      }
+      var = var->next;
+    }
+    if (res == 1 && !PyErr_Occurred()) {
+      PyErr_SetString(PyExc_NameError,"Unknown C global variable");
+    }
+    return res;
+  }
+  
+  SWIGINTERN PyTypeObject*
+  swig_varlink_type(void) {
+    static char varlink__doc__[] = "Swig var link object";
+    static PyTypeObject varlink_type;
+    static int type_init = 0;
+    if (!type_init) {
+      const PyTypeObject tmp = {
+        /* PyObject header changed in Python 3 */
+#if PY_VERSION_HEX >= 0x03000000
+        PyVarObject_HEAD_INIT(NULL, 0)
+#else
+        PyObject_HEAD_INIT(NULL)
+        0,                                  /* ob_size */
+#endif
+        (char *)"swigvarlink",              /* tp_name */
+        sizeof(swig_varlinkobject),         /* tp_basicsize */
+        0,                                  /* tp_itemsize */
+        (destructor) swig_varlink_dealloc,  /* tp_dealloc */
+        (printfunc) swig_varlink_print,     /* tp_print */
+        (getattrfunc) swig_varlink_getattr, /* tp_getattr */
+        (setattrfunc) swig_varlink_setattr, /* tp_setattr */
+        0,                                  /* tp_compare */
+        (reprfunc) swig_varlink_repr,       /* tp_repr */
+        0,                                  /* tp_as_number */
+        0,                                  /* tp_as_sequence */
+        0,                                  /* tp_as_mapping */
+        0,                                  /* tp_hash */
+        0,                                  /* tp_call */
+        (reprfunc) swig_varlink_str,        /* tp_str */
+        0,                                  /* tp_getattro */
+        0,                                  /* tp_setattro */
+        0,                                  /* tp_as_buffer */
+        0,                                  /* tp_flags */
+        varlink__doc__,                     /* tp_doc */
+        0,                                  /* tp_traverse */
+        0,                                  /* tp_clear */
+        0,                                  /* tp_richcompare */
+        0,                                  /* tp_weaklistoffset */
+#if PY_VERSION_HEX >= 0x02020000
+        0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0,0, /* tp_iter -> tp_weaklist */
+#endif
+#if PY_VERSION_HEX >= 0x02030000
+        0,                                  /* tp_del */
+#endif
+#if PY_VERSION_HEX >= 0x02060000
+        0,                                  /* tp_version */
+#endif
+#ifdef COUNT_ALLOCS
+        0,0,0,0                             /* tp_alloc -> tp_next */
+#endif
+      };
+      varlink_type = tmp;
+      type_init = 1;
+#if PY_VERSION_HEX < 0x02020000
+      varlink_type.ob_type = &PyType_Type;
+#else
+      if (PyType_Ready(&varlink_type) < 0)
+      return NULL;
+#endif
+    }
+    return &varlink_type;
+  }
+  
+  /* Create a variable linking object for use later */
+  SWIGINTERN PyObject *
+  SWIG_Python_newvarlink(void) {
+    swig_varlinkobject *result = PyObject_NEW(swig_varlinkobject, swig_varlink_type());
+    if (result) {
+      result->vars = 0;
+    }
+    return ((PyObject*) result);
+  }
+  
+  SWIGINTERN void 
+  SWIG_Python_addvarlink(PyObject *p, char *name, PyObject *(*get_attr)(void), int (*set_attr)(PyObject *p)) {
+    swig_varlinkobject *v = (swig_varlinkobject *) p;
+    swig_globalvar *gv = (swig_globalvar *) malloc(sizeof(swig_globalvar));
+    if (gv) {
+      size_t size = strlen(name)+1;
+      gv->name = (char *)malloc(size);
+      if (gv->name) {
+        strncpy(gv->name,name,size);
+        gv->get_attr = get_attr;
+        gv->set_attr = set_attr;
+        gv->next = v->vars;
+      }
+    }
+    v->vars = gv;
+  }
+  
+  SWIGINTERN PyObject *
+  SWIG_globals(void) {
+    static PyObject *_SWIG_globals = 0; 
+    if (!_SWIG_globals) _SWIG_globals = SWIG_newvarlink();  
+    return _SWIG_globals;
+  }
+  
+  /* -----------------------------------------------------------------------------
+   * constants/methods manipulation
+   * ----------------------------------------------------------------------------- */
+  
+  /* Install Constants */
+  SWIGINTERN void
+  SWIG_Python_InstallConstants(PyObject *d, swig_const_info constants[]) {
+    PyObject *obj = 0;
+    size_t i;
+    for (i = 0; constants[i].type; ++i) {
+      switch(constants[i].type) {
+      case SWIG_PY_POINTER:
+        obj = SWIG_InternalNewPointerObj(constants[i].pvalue, *(constants[i]).ptype,0);
+        break;
+      case SWIG_PY_BINARY:
+        obj = SWIG_NewPackedObj(constants[i].pvalue, constants[i].lvalue, *(constants[i].ptype));
+        break;
+      default:
+        obj = 0;
+        break;
+      }
+      if (obj) {
+        PyDict_SetItemString(d, constants[i].name, obj);
+        Py_DECREF(obj);
+      }
+    }
+  }
+  
+  /* -----------------------------------------------------------------------------*/
+  /* Fix SwigMethods to carry the callback ptrs when needed */
+  /* -----------------------------------------------------------------------------*/
+  
+  SWIGINTERN void
+  SWIG_Python_FixMethods(PyMethodDef *methods,
+    swig_const_info *const_table,
+    swig_type_info **types,
+    swig_type_info **types_initial) {
+    size_t i;
+    for (i = 0; methods[i].ml_name; ++i) {
+      const char *c = methods[i].ml_doc;
+      if (c && (c = strstr(c, "swig_ptr: "))) {
+        int j;
+        swig_const_info *ci = 0;
+        const char *name = c + 10;
+        for (j = 0; const_table[j].type; ++j) {
+          if (strncmp(const_table[j].name, name, 
+              strlen(const_table[j].name)) == 0) {
+            ci = &(const_table[j]);
+            break;
+          }
+        }
+        if (ci) {
+          void *ptr = (ci->type == SWIG_PY_POINTER) ? ci->pvalue : 0;
+          if (ptr) {
+            size_t shift = (ci->ptype) - types;
+            swig_type_info *ty = types_initial[shift];
+            size_t ldoc = (c - methods[i].ml_doc);
+            size_t lptr = strlen(ty->name)+2*sizeof(void*)+2;
+            char *ndoc = (char*)malloc(ldoc + lptr + 10);
+            if (ndoc) {
+              char *buff = ndoc;
+              strncpy(buff, methods[i].ml_doc, ldoc);
+              buff += ldoc;
+              strncpy(buff, "swig_ptr: ", 10);
+              buff += 10;
+              SWIG_PackVoidPtr(buff, ptr, ty->name, lptr);
+              methods[i].ml_doc = ndoc;
+            }
+          }
+        }
+      }
+    }
+  } 
+  
+#ifdef __cplusplus
+}
+#endif
+
+/* -----------------------------------------------------------------------------*
+ *  Partial Init method
+ * -----------------------------------------------------------------------------*/
+
+#ifdef __cplusplus
+extern "C"
+#endif
+
+SWIGEXPORT 
+#if PY_VERSION_HEX >= 0x03000000
+PyObject*
+#else
+void
+#endif
+SWIG_init(void) {
+  PyObject *m, *d, *md;
+#if PY_VERSION_HEX >= 0x03000000
+  static struct PyModuleDef SWIG_module = {
+# if PY_VERSION_HEX >= 0x03020000
+    PyModuleDef_HEAD_INIT,
+# else
+    {
+      PyObject_HEAD_INIT(NULL)
+      NULL, /* m_init */
+      0,    /* m_index */
+      NULL, /* m_copy */
+    },
+# endif
+    (char *) SWIG_name,
+    NULL,
+    -1,
+    SwigMethods,
+    NULL,
+    NULL,
+    NULL,
+    NULL
+  };
+#endif
+  
+#if defined(SWIGPYTHON_BUILTIN)
+  static SwigPyClientData SwigPyObject_clientdata = {
+    0, 0, 0, 0, 0, 0, 0
+  };
+  static PyGetSetDef this_getset_def = {
+    (char *)"this", &SwigPyBuiltin_ThisClosure, NULL, NULL, NULL
+  };
+  static SwigPyGetSet thisown_getset_closure = {
+    (PyCFunction) SwigPyObject_own,
+    (PyCFunction) SwigPyObject_own
+  };
+  static PyGetSetDef thisown_getset_def = {
+    (char *)"thisown", SwigPyBuiltin_GetterClosure, SwigPyBuiltin_SetterClosure, NULL, &thisown_getset_closure
+  };
+  PyObject *metatype_args;
+  PyTypeObject *builtin_pytype;
+  int builtin_base_count;
+  swig_type_info *builtin_basetype;
+  PyObject *tuple;
+  PyGetSetDescrObject *static_getset;
+  PyTypeObject *metatype;
+  SwigPyClientData *cd;
+  PyObject *public_interface, *public_symbol;
+  PyObject *this_descr;
+  PyObject *thisown_descr;
+  int i;
+  
+  (void)builtin_pytype;
+  (void)builtin_base_count;
+  (void)builtin_basetype;
+  (void)tuple;
+  (void)static_getset;
+  
+  /* metatype is used to implement static member variables. */
+  metatype_args = Py_BuildValue("(s(O){})", "SwigPyObjectType", &PyType_Type);
+  assert(metatype_args);
+  metatype = (PyTypeObject *) PyType_Type.tp_call((PyObject *) &PyType_Type, metatype_args, NULL);
+  assert(metatype);
+  Py_DECREF(metatype_args);
+  metatype->tp_setattro = (setattrofunc) &SwigPyObjectType_setattro;
+  assert(PyType_Ready(metatype) >= 0);
+#endif
+  
+  /* Fix SwigMethods to carry the callback ptrs when needed */
+  SWIG_Python_FixMethods(SwigMethods, swig_const_table, swig_types, swig_type_initial);
+  
+#if PY_VERSION_HEX >= 0x03000000
+  m = PyModule_Create(&SWIG_module);
+#else
+  m = Py_InitModule((char *) SWIG_name, SwigMethods);
+#endif
+  md = d = PyModule_GetDict(m);
+  
+  SWIG_InitializeModule(0);
+  
+#ifdef SWIGPYTHON_BUILTIN
+  SwigPyObject_stype = SWIG_MangledTypeQuery("_p_SwigPyObject");
+  assert(SwigPyObject_stype);
+  cd = (SwigPyClientData*) SwigPyObject_stype->clientdata;
+  if (!cd) {
+    SwigPyObject_stype->clientdata = &SwigPyObject_clientdata;
+    SwigPyObject_clientdata.pytype = SwigPyObject_TypeOnce();
+  } else if (SwigPyObject_TypeOnce()->tp_basicsize != cd->pytype->tp_basicsize) {
+    PyErr_SetString(PyExc_RuntimeError, "Import error: attempted to load two incompatible swig-generated modules.");
+# if PY_VERSION_HEX >= 0x03000000
+    return NULL;
+# else
+    return;
+# endif
+  }
+  
+  /* All objects have a 'this' attribute */
+  this_descr = PyDescr_NewGetSet(SwigPyObject_type(), &this_getset_def);
+  (void)this_descr;
+  
+  /* All objects have a 'thisown' attribute */
+  thisown_descr = PyDescr_NewGetSet(SwigPyObject_type(), &thisown_getset_def);
+  (void)thisown_descr;
+  
+  public_interface = PyList_New(0);
+  public_symbol = 0;
+  (void)public_symbol;
+  
+  PyDict_SetItemString(md, "__all__", public_interface);
+  Py_DECREF(public_interface);
+  for (i = 0; SwigMethods[i].ml_name != NULL; ++i)
+  SwigPyBuiltin_AddPublicSymbol(public_interface, SwigMethods[i].ml_name);
+  for (i = 0; swig_const_table[i].name != 0; ++i)
+  SwigPyBuiltin_AddPublicSymbol(public_interface, swig_const_table[i].name);
+#endif
+  
+  SWIG_InstallConstants(d,swig_const_table);
+  
+#if PY_VERSION_HEX >= 0x03000000
+  return m;
+#else
+  return;
+#endif
+}
+
Index: Pi/C/examples/capi/setup.py.orig
===================================================================
--- Pi/C/examples/capi/setup.py.orig	(revision 0)
+++ Pi/C/examples/capi/setup.py.orig	(revision 7)
@@ -0,0 +1,8 @@
+from distutils.core import setup, Extension
+
+module1 = Extension('fib', sources = ['fibmodule.c'])
+
+setup (name = 'PackageName',
+       version = '1.0',
+       description = 'This is a demo package',
+       ext_modules = [module1])
Index: Pi/C/examples/capi/fibmodule.c
===================================================================
--- Pi/C/examples/capi/fibmodule.c	(revision 0)
+++ Pi/C/examples/capi/fibmodule.c	(revision 7)
@@ -0,0 +1,30 @@
+#include <Python.h>
+ 
+int _fib(int n)
+{
+  if (n < 2)
+    return n;
+  else
+    return _fib(n-1) + _fib(n-2);
+}
+ 
+static PyObject* fib(PyObject* self, PyObject* args)
+{
+  int n;
+ 
+  if (!PyArg_ParseTuple(args, "i", &n))
+    return NULL;
+ 
+  return Py_BuildValue("i", _fib(n));
+}
+ 
+static PyMethodDef FibMethods[] = {
+  {"fib", fib, METH_VARARGS, "Calculate the Fibonacci numbers."},
+  {NULL, NULL, 0, NULL}
+};
+ 
+PyMODINIT_FUNC
+initfib(void)
+{
+  (void) Py_InitModule("fib", FibMethods);
+}
Index: Pi/MPI/PiMPI.py
===================================================================
--- Pi/MPI/PiMPI.py	(revision 0)
+++ Pi/MPI/PiMPI.py	(revision 7)
@@ -0,0 +1,79 @@
+#!/usr/bin/env python
+#
+# PiMC model using mpi4py MPI implementation for Python
+#
+# CC BY-NC-SA 2011 : <emmanuel.quemener@ens-lyon.fr> 
+#
+# Thanks to Lisandro Dalcin for MPI4PY :
+# http://mpi4py.scipy.org/
+
+import sys
+from math import exp
+from random import random
+import time
+# MPI librairie call
+import mpi4py
+from mpi4py import MPI
+
+def MainLoop(iterations):
+
+    total=0
+    for i in xrange(iterations):
+        # Random access coordonate
+        x,y=random(),random()
+
+	if ((x*x+y*y) < 1.0):
+		total+=1
+
+    return(total)
+
+if __name__=='__main__':
+	
+   # MPI Init
+   comm = MPI.COMM_WORLD
+   rank = comm.Get_rank()
+	
+   # Define number of Nodes on with computing is performed (exclude 0)
+   NODES=comm.Get_size()
+
+   # Au dessus de 4 10^7, le MPI ne veut plus se lancer...
+   Iterations=1000000000
+
+   # pass explicit MPI datatypes
+   total=0
+   if rank == 0:
+	   # Define iterations to send to each node
+	   if Iterations%NODES==0:
+		   iterations=Iterations/NODES
+	   else:
+		   iterations=Iterations/NODES+1
+	   print "%i iterations will be send to each node" % iterations
+	   
+	   for i in range(1,NODES):
+		   
+		   print "Send from 0 to node %i" % i
+		   ToSend=iterations
+		   # Send MPI call to each node
+		   comm.send(ToSend, dest=i, tag=11)
+
+	   # Master does part of job !
+	   total=MainLoop(iterations)
+	   print "Partial Result from master %i: %i" % (rank,total)
+
+	   print "Retreive results..."
+	   for i in range(1,NODES):
+		   # Receive MPI call from each node
+		   Output=comm.recv(source=i,tag=11)
+		   print "Partial Result from %i: %s" % (i,Output)
+		   total+=Output
+
+	   print "Global Result: %i" % (total)
+   else:
+	   # Slave applies simulation to set provided by master
+	   # Receive MPI call with Input set
+	   ToReceive=comm.recv(source=0, tag=11)
+	   iterations=ToReceive
+	   print "Rank %i receives with job with %i" % (rank,iterations) 
+	   Output=MainLoop(iterations)
+
+	   comm.send(Output, dest=0, tag=11)

Property changes on: Pi/MPI/PiMPI.py
___________________________________________________________________
Added: svn:executable
   + *

Index: Pi/MPI/PostProcess.py
===================================================================
--- Pi/MPI/PostProcess.py	(revision 0)
+++ Pi/MPI/PostProcess.py	(revision 7)
@@ -0,0 +1,9 @@
+import numpy
+npzfile=numpy.load("Ising2D_MPI_256_16777216.npz")
+npzfile['arr_0']
+J=npzfile['arr_0'].transpose()[0]
+E=npzfile['arr_0'].transpose()[1]
+M=npzfile['arr_0'].transpose()[2]
+import matplotlib.pyplot 
+matplotlib.pyplot.plot(J,E)
+matplotlib.pyplot.show()
Index: Pi/RNG/MWC.py
===================================================================
--- Pi/RNG/MWC.py	(revision 0)
+++ Pi/RNG/MWC.py	(revision 7)
@@ -0,0 +1,22 @@
+#!/usr/bin/env python
+
+import numpy
+
+SIZE=256
+
+if __name__=='__main__':
+
+    z=numpy.uint32(37)
+    w=numpy.uint32(91)
+
+    for i in range(1000):
+
+        z=numpy.uint32(36969*(z&65535)+(z>>16))
+        w=numpy.uint32(18000*(w&65535)+(w>>16))
+        MWC= lambda : numpy.uint32((z<<16)+w)
+        MWCfp= lambda: (MWC() + 1.0) * 2.328306435454494e-10
+
+        print i,MWC(),numpy.uint32(MWCfp()*SIZE)
+
+
+

Property changes on: Pi/RNG/MWC.py
___________________________________________________________________
Added: svn:executable
   + *

Index: Pi/RNG/MWC.c
===================================================================
--- Pi/RNG/MWC.c	(revision 0)
+++ Pi/RNG/MWC.c	(revision 7)
@@ -0,0 +1,21 @@
+#include <stdio.h>
+#include <math.h>
+
+#define znew (z=36969*(z&65535)+(z>>16))
+#define wnew (w=18000*(w&65535)+(w>>16))
+#define MWC ((znew<<16)+wnew)
+#define MWCfp (MWC + 1.0f) * 2.328306435454494e-10f
+
+int main(void) {
+  int i,z,w;
+
+  z=37;
+  w=91;
+
+  for (i=1;i<100;i++) {
+
+    printf("%i %i %i %i\n",i,znew,wnew,MWC);
+  
+  }
+
+}
