Programación y computación paralelas

   #pragma omp nombre_directiva [cláusulas, ...] 

• #pragma omp requiere a todas las directivas OpenMP para C/C++.

• nombre_directiva es un nombre válido de directiva, y tiene que aparecer después de pragma y antes de cualquier cláusula. En nuestro ejemplo, es parallel for.

• [cláusulas, ...] son opcionales. Las cláusulas pueden ir en cualquier orden y repetirse cuando sea necesario, a menos que haya alguna restricción. En nuestro caso, son reduction y private.

   #pragma omp parallel [cláusulas]
      bloque

• Se crea un grupo de flujos y el flujo que los pone en marcha toma el rol de flujo maestro.

• El número de flujos que hay que crear se obtiene por la variable de entorno OMP_NUM_THHREADS o, de manera explícita, con una llamada a la librería, tal y como veremos más adelante. Si el valor es igual a cero, se ejecuta de manera secuencial.

• Hay una barrera implícita al final de la región, de modo que el flujo maestro espera a que acaben todos los esclavos para continuar con la ejecución secuencial.

• Cuando dentro de una región hay otro constructor parallel, cada esclavo crea otro grupo de flujos esclavos de los que sería el maestro. Esto se denomina paralelismo imbricado y, pese a que se pueden programar de esta manera, en algunas implementaciones de OpenMP la creación de grupos de esclavos dentro de una región paralela no está soportada.

• Las cláusulas de compartición de variables que soporta la directiva parallel son principalmente private, firstprivate, lastprivate, default, shared, copyin y reduction. El significado de estas cláusulas, así como el de otras de compartición, se verá más adelante.

• private(lista): las variables de la lista son privadas a los flujos, lo que quiere decir que cada flujo tiene una variable privada con este nombre (que pueden tener valores diferentes en distintos flujos). Las variables no se inicializan antes de entrar en la región paralela (por lo tanto, no podemos esperar un valor inicial concreto al entrar en la región paralela), y no se guarda su valor al salir (de modo que no podemos esperar que el último valor de una variable privada se conserve al finalizar la región paralela).

• firstprivate(lista): las variables son privadas a los flujos y se inicializan al entrar en la región paralela con el valor que tuviera la variable correspondiente del flujo maestro.

• lastprivate(lista): son privadas a los flujos y al salir de la región paralela quedan con el valor de la última iteración (si estamos en un bucle for paralelo) o sección (veremos más adelante el funcionamiento de las secciones).

• shared(lista): indica las variables compartidas por todos los flujos. Por defecto todas son compartidas, por lo que no es realmente necesario utilizar esta cláusula en la mayoría de los casos.

• default(shared|none): indica cómo serán las variables por defecto. Si se especifica none, habrá que señalar de manera explícita con la cláusula shared las que se desea que sean compartidas.

• reduction(operador:lista): las variables de la lista se obtienen por la aplicación del operador, que tiene que ser asociativo.

• copyin(lista): se utiliza para asignar el valor de la variable en el maestro a variables del tipo threadprivate.

   #pragma omp for [cláusulas]
      bucle for

• Las iteraciones se ejecutan en paralelo por los flujos que ya existen, creados de manera previa con parallel.

• El bucle for debe tener una forma especial (forma canónica), tal y como muestra la figura 9. La parte de inicialización ha de tener una asignación; la parte del incremento, una suma o resta; la de de evaluación es la comparación de una variable entera con signo con un valor, utilizando un comparador mayor o menor (puede incluir igual); y los valores que aparecen en las tres partes de for deben ser enteros.

Figura 9. Forma canónica de los bucles en OpenMP

• Hay una barrera implícita al final del bucle, a no ser que se utilice la cláusula nowait.

• Las cláusulas de compartición de variables que admite son private, firstprivate, lastprivate y reduction.

• Puede aparecer una cláusula schedule para indicar de qué manera se dividen las iteraciones de for entre los flujos, es decir, la política de planificación.

• schedule(static, tamaño): las iteraciones se dividen según el tamaño de bloque que se indica. Por ejemplo, si el número de iteraciones del bucle es 100, están numeradas de 0 a 99 y el tamaño de bloque es 2, se consideran 50 bloques, cada uno de estos con dos iteraciones (bloques con iteraciones 0-1, 2-3, etc.). Si disponemos de 4 flujos, los bloques de iteraciones se asignan a los flujos de manera cíclica, con lo que al flujo 0 le corresponden las iteraciones 0-1, 8-9, etc.; al flujo 1, las iteraciones 2-3, 10-11, etc.; al 2, las 4-5, 12-13, etc.; y al 3, las 6-7, 14-15, etc. Si no se indica ningún tamaño de bloque, las iteraciones se dividen por igual entre los flujos y los bloques son de tamaño máximo: si tenemos 12 iteraciones y 4 flujos, se asignan al flujo 0 las iteraciones 0-2, al 1 las 3-5, al 2 las 6-8 y al 4, las 9-11.

• schedule(dynamic, tamaño): las iteraciones se agrupan según el tamaño de bloque especificado y se asignan a los flujos de manera dinámica cuando van acabando su trabajo. En el caso anterior de 100 iteraciones, tamaño de bloque 2 y 4 flujos, se asignan inicialmente al flujo i las iteraciones 2i y 2i + 1, y a partir de aquí el resto de los bloques de dos iteraciones se asignan a los flujos según se vayan quedando sin trabajo. Cuando el volumen de computación de cada iteración no se conoce a priori, quizá sea preferible utilizar una asignación dinámica, puesto que permite reducir el desbalanceo.

• schedule(guided, tamaño): las iteraciones se asignan dinámicamente a los flujos pero con tamaño de bloque decreciente hasta llegar al tamaño de bloque que se indica (si no se indica ningún valor, es 1 por defecto). El tamaño inicial de bloque y la forma en que decrece dependen de su implementación.

• schedule(runtime): decide la política de planificación en tiempo de ejecución mediante la variable de entorno OMP_SCHEDULE.

Figura 10. Trazas de ejecución extraídas con la aplicación Paraver del Barcelona Supercomputing Center

   #pragma omp sections [cláusulas]
      {
         [#pragma omp section]
            bloque
         [#pragma omp section
            bloque
         ...
         ]
      }

• Cada sección se ejecuta mediante un flujo. La forma en que se distribuyen las secciones entre los flujos depende de la implementación concreta de OpenMP.

• Hay una barrera final, a menos que se utilice la cláusula nowait.

• Las cláusulas de compartición de variables que admite son private, firstprivate, lastprivate y reduction .

   #pragma omp parallel for [cláusulas]
      bucle for

   #pragma omp parallel sections [cláusulas]

• single: el código afectado por la directiva lo ejecutará un único flujo. Los flujos que no están trabajando durante la ejecución de la directiva esperan al final. Admite las cláusulas private, firstprivate y nowait. No está permitido hacer bifurcaciones hacia/desde un bloque single. Resulta útil para secciones de código que pueden ser no seguras para su ejecución paralela (por ejemplo, para entrada/salida).

• master: el código lo ejecuta solo el flujo maestro. El resto de los flujos no ejecutan esta sección de código.

• critical: protege una sección de código para que pueda acceder un solo flujo al mismo tiempo. Se le puede asociar un nombre de la forma:

              #pragma omp critical [nombre]

  Por lo que puede haber secciones críticas protegiendo zonas diferentes del programa, de modo que varios flujos a la vez pueden acceder a secciones con nombres diferentes. Las regiones que tengan el mismo nombre se tratan como la misma región. Todas las que no tienen nombre se consideran la misma. No está permitido hacer bifurcaciones hacia/desde un bloque critical.

• atomic: asegura que una posición de memoria se modifique sin que múltiples flujos intenten escribirla de manera simultánea. Se aplica a la sentencia que sigue a la directiva. Solo se asegura en modo exclusivo la actualización de la variable, pero no la evaluación de la expresión.

• barrier: sincroniza todos los flujos. Cuando un flujo llega a la barrera espera a que lleguen los otros y, cuando han llegado todos, siguen su ejecución.

• threadprivate: se utiliza para que variables globales se conviertan en locales y persistentes a un flujo a través de múltiples regiones paralelas.

• ordered: asegura que el código se ejecute en el orden en el que las iteraciones se ejecutan en su forma secuencial. Puede aparecer solo una vez en el contexto de una directiva for o parallel for. No está permitido hacer bifurcaciones hacia/desde un bloque ordered. Solo puede haber un único flujo ejecutándose de manera simultánea en una sección ordered. Una iteración de un bucle no puede ejecutar la misma directiva ordered más de una vez, y no tiene que ejecutar más de una directiva ordered. Un bucle con una directiva ordered debe contener una cláusula ordered.

• flush: tiene la forma siguiente.

              #pragma omp flush [lista-de-variables]

  Asegura que el valor de las variables se actualiza en todos los flujos en los que son visibles. Si no hay lista de variables, se actualizarán todas.

• omp_get_max_threads: obtiene la máxima cantidad posible de flujos.

• omp_get_num_procs: retorna el número máximo de procesadores que se pueden asignar al programa.

• omp_in_parallel: retorna un valor diferente a cero si se ejecuta dentro de una región paralela.

• void omp_init_lock(omp_lock_t *lock): para inicializar un lock, que se inicializa como no bloqueado.

• void omp_init_destroy(omp_lock_t *lock): para destruir un lock.

• void omp_set_lock(omp_lock_t *lock): para bloquear un lock.

• void omp_unset_lock(omp_lock_t *lock): para desbloquear un lock.

• void omp_test_lock(omp_lock_t *lock): para comprobar si un lock está bloqueado o no y así evitar bloqueos indeseados.

   gcc -fopenmp -o programa programa.c 

   export OMP_NUM_THREADS=6

• OMP_SCHEDULE indica el tipo de planificación para for y parallel for.

• OMP_DYNAMIC autoriza o desautoriza el ajuste dinámico del número de flujos.

• OMP_NESTED autoriza o desautoriza el paralelismo imbricado. Por defecto, no está autorizado.

• Reservar memoria en el dispositivo (paso 1 de la figura 15).

• Transferir los datos necesarios desde el host al espacio de memoria asignado al dispositivo (paso 2 de la figura 15).

• Invocar la ejecución del kernel en cuestión (paso 3 de la figura 15).

• Transferir los datos con los resultados desde el dispositivo hacia el host y liberar la memoria del dispositivo (si ya no es necesaria), una vez finalizada la ejecución del kernel (paso 4 de la figura 15).

Figura 15. Esquema de los pasos para la ejecución de un kernel en una GPU

void matrix_add_cpu (fload *A, float *B, float *C, int N)
{
   int i, j, index;
   for (i=0; i<N; i++){
      for (j=0; j<N; j++){
         index = i+j*N;
         C[index] = A[index] + B[index];
      }
   } 
}

int main(){
   matrix_add_cpu(a, b, c, N);
}

• Acceso de lectura/escritura a la memoria global, por grid.

• Acceso solo de lectura a la memoria constante, por grid.

• Acceso de lectura/escritura a los registros, por flujo.

• Acceso de lectura/escritura a la memoria local, por flujo.

• Acceso de lectura/escritura a la memoria compartida, por bloque.

• Acceso de lectura/escritura.

Figura 16. Modelo de memoria de CUDA

   float *Matriz;
   int tamaño = ANCHURA * LONGITUD * sizeof(float);
   cudaMalloc((void **) &Matriz, tamaño);
   ...
   cudaFree(Matriz);

• cudaMemcpyHostToHost: de la memoria del host hacia la memoria del mismo host.

• cudaMemcpyHostToDevice: de la memoria del host hacia la memoria del dispositivo.

• cudaMemcpyDeviceToHost: de la memoria del dispositivo hacia la memoria del host.

• cudaMemcpyDeviceToDevice: de la memoria del dispositivo hacia la memoria del dispositivo.

void matrix_add (fload *A, float *B, float *C, int N)
{
   int size = N * N * sizeof(float);
   float * Ad, Bd, Cd;

   // 1. Reservar memoria para las matrices
   cudaMalloc((void **) &Ad, size);
   cudaMalloc((void **) &Bd, size);
   cudaMalloc((void **) &Cd, size);
   // 2. Copiar matrices de entrada al dispositivo
   cudaMemcpy(Ad, A, size, cudaMemcpyHostToDevice);
   cudaMemcpy(Bd, B, size, cudaMemcpyHostToDevice);
   ...
   // 4. Copiar resultado C hacia el host
   cudaMemcpy(C, Cd, size, cudaMemcpyDeviceToHost);
   ...

__global__ matrix_add_gpu (fload *A, float *B, float *C, int N)
{
   int i = blockIdx.x * blockDim.x + threadIdx.x;
   int j = blockIdx.y * blockDim.y + threadIdx.y;
   int index = i + j*N;
   if (i<N && j<N){
         C[index] = A[index] + B[index];
   }
}

int main(){
   dim3 dimBlock(blocksize, blocksize);
   dim3 dimGrid(N/dimBlock.x, N/dimBlock.y);
   matrix_add_gpu<<<dimGrid, dimBlock>>>(a, b, c, N);
}

• La palabra clave __device__ indica que la función declarada es una función CUDA de dispositivo. Una función de dispositivo se ejecuta únicamente en un dispositivo CUDA y solo se puede llamar desde un kernel o desde otra función de dispositivo. Estas funciones no pueden tener ni llamadas recursivas ni llamadas indirectas a funciones mediante punteros.

• La palabra clave __host__ indica que la función es una función de host, es decir, una función simple de C que se ejecuta en el host y, por lo tanto, que puede ser llamada desde cualquier función de host. Por defecto, todas las funciones en un programa CUDA son funciones de host si no se especifica ninguna palabra clave en la definición de la función.

   // Configuración de las dimensiones de grid y bloques
   dim3 dimGrid(2, 2, 1);
   dim3 dimBlock(4, 2, 2);

   // Invocación del kernel (suma de matrices)
   matrix_add_gpu<<<dimGrid, dimBlock>>>(a, b, c, N);

• La memoria global es la que pueden utilizar todas las unidades de cálculo de un dispositivo.

• La memoria constante es la memoria que todas las unidades de cálculo de un dispositivo pueden utilizar para almacenar datos constantes para acceso solo de lectura durante la ejecución de un kernel. El procesador host es responsable de asignar e iniciar los objetos de memoria que residen en el espacio de memoria.

• La memoria local es la memoria que pueden utilizar los trabajos elementales de un grupo.

• La memoria privada es la que solo puede utilizar una unidad de cálculo. Esto es similar a los registros en una sola unidad de cálculo o un solo núcleo de una CPU.

Figura 21. Arquitectura y jerarquía de memoria de un dispositivo OpenCL

__kernel void matrix_add_opencl (   __global const float *A, 
            __global const float *B, 
            __global float *C,
            int N) {
   int i = get_global_id(0);
   int j = get_global_id(1);
   int index = i + j*N;
   if (i<N && j<N){
         C[index] = A[index] + B[index];
   }
}

main(){
   // Inicialización de variables, etc.
   (...)

   // 1. Creación del contexto y cola en el dispositivo
   cl_context context = clCreateContextFromType(0, CL_DEVICE_TYPE_GPU, NULL, NULL, NULL); 
   // Para obtener la lista de dispositivos GPU asociados al contexto 
   size_t cb;
   clGetContextInfo( context,   CL_CONTEXT_DEVICES,   0,   NULL, &cb);
   cl_device_id *devices = malloc(cb);
   clGetContextInfo( context,   CL_CONTEXT_DEVICES,   cb,   devices,   NULL);
   cl_cmd_queue cmd_queue = clCreateCommandQueue(context, devices[0], 0 , NULL); 

   // 2. Definición de los objetos en memoria (matrices A, B y C)
   cl_mem memobjs[3];
   memobjs[0] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
            sizeof(cl_float)*n, srcA, NULL);
   memobjs[1] = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, 
            sizeof(cl_float)*n, srcB, NULL);
   memobjs[2] = clCreateBuffer(context, CL_MEM_WRITE_ONLY, sizeof(cl_float)*n, NULL, NULL);

   // 3. Definición del kernel y argumentos
   cl_program program = clCreateProgramWithSource(context, 1, &program_source, NULL, NULL);
   cl_int err = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
   cl_kernel kernel = clCreateKernel(program, "matrix_add_opencl", NULL);
   err = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&memobjs[0]);
   err |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&memobjs[1]);
   err |= clSetKernelArg(kernel, 2, sizeof(cl_mem), (void *)&memobjs[2]);
   err |= clSetKernelArg(kernel, 3, sizeof(int), (void *)&N);

   // 4. Invocación del kernel
   size_t global_work_size[1] = n;
   err = clEnqueueNDRangeKernel(cmd_queue, kernel, 1, NULL, global_work_size, 
      NULL, 0, NULL, NULL);

   // 5. Lectura de los resultados (matriz C)
   err = clEnqueueReadBuffer(context, memobjs[2], CL_TRUE, 0, n*sizeof(cl_float), 
      dstC, 0, NULL, NULL);
   (...)

int MPI_Send(void *buf, int count, MPI_Datatype datatype, 
                  int dest, int tag, MPI_Comm comm)

int MPI_Recv(void *buf, int count, MPI_Datatype datatype,
                  int source, int tag, MPI_Comm comm, 
                  MPI_Status *status)

• buf contiene el inicio de la zona de memoria de la que se tomarán los datos que hay que enviar o bien donde se almacenarán los datos que se reciben.

• count indica el número de datos que hay que enviar o el espacio disponible para recibir (no el número de datos del mensaje que se recibe, puesto que el tamaño lo determina quien envía).

• datatype es el tipo de datos que hay que transferir y debe ser un tipo MPI (MPI_Datatype), por ejemplo MPI_CHAR o MPI_INT.

• dest y source son los identificadores del proceso al que se envía el mensaje y del proceso del que se recibe, de manera respectiva. Se puede utilizar la constante MPI_ANY_SOURCE para indicar que es posible recibirlo desde cualquier proceso.

• tag se utiliza para diferenciar entre mensajes, y su valor debe coincidir en el proceso que envía y el que recibe. Se puede utilizar MPI_ANY_TAG para indicar que el mensaje es compatible con mensajes con cualquier identificador.

• comm es el comunicador dentro del que se hace la comunicación. Es del tipo MPI_Comm, y en el ejemplo se utiliza el identificador de comunicador formado por todos los procesos (MPI_COMM_WORLD).

• status referencia una variable de tipo MPI_Status. En el programa no se utiliza, pero contiene información del mensaje que se ha recibido y se puede consultar para identificar alguna característica del mensaje. Por ejemplo, su longitud, el proceso de origen, etc.

   MPI_Comm_rank(MPI_COMM_WORLD,&myrank); /* find rank */
   if (myrank == 0) {
         int x;
         MPI_Send(&x, 1, MPI_INT, 1, msgtag, MPI_COMM_WORLD);
   } 
   else if (myrank == 1) {
         int x;
         MPI_Recv(&x, 1, MPI_INT, 0, msgtag, MPI_COMM_WORLD,status);
}

int MPI_Isend(void *buf, int count, MPI_Datatype datatype, 
         int dest, int tag, MPI_Comm comm)

int MPI_Irecv(void *buf, int count, MPI_Datatype datatype,
      int source, int tag, MPI_Comm comm, 
      MPI_Status *status)

   MPI_Comm_rank(MPI_COMM_WORLD, &myrank); /* find rank */
   if (myrank == 0) {
      int x;
      MPI_Isend(&x,1,MPI_INT, 1, msgtag, MPI_COMM_WORLD, req1);
      compute(); // Efectúa algún cálculo mientras se hace el envío
      MPI_Wait(req1, status);
   } 
   else if (myrank == 1) {
      int x;
      MPI_Recv(&x,1,MPI_INT,0,msgtag, MPI_COMM_WORLD, status);
}

int MPI_Barrier(MPI_Comm comm) 

int MPI_Bcast(void *buffer, int count, 
      MPI_Datatype datatype, int root, MPI_comm comm) 

int MPI_Scatter (void *sendbuf, int sendcount, 
            MPI_Datatype sendtype, void *recvbuf, 
            int recvcount, MPI_Datatype recvtype, 
            int root, MPI_Comm comm)

int MPI_Gather(void *sendbuf, int sendcount, 
            MPI_Datatype sendtype, void *recvbuf, 
            int recvcount, MPI_Datatype recvtype, 
            int root, MPI_Comm comm)

int MPI_Allgather (void *sendbuf, int sendcount, 
            MPI_Datatype sendtype, void *recvbuf, 
            int recvcount, MPI_Datatype recvtype, 
            MPI_Comm comm)

int MPI_Alltoall(void *sendbuf, int sendcount, 
            MPI_Datatype sendtype, void *recvbuf, 
            int recvcount, MPI_Datatype recvtype, 
            MPI_Comm comm)

int MPI_Reduce(void *sendbuf, void *recvbuf, int count, 
            MPI_Datatype datatype, MPI_Op op, 
            int root, MPI_Comm comm)

#include "mpi.h"
#include <stdio.h>
int main(int argc, char **argv)
{
   int MyProc, tag=1, size;
   char msg='A', msg_recpt ;
   MPI_Status *status ;
   int root ;
   int left, right, interval ;
   int number, start, end, sum, GrandTotal;
   int mystart, myend;
   MPI_Init(&argc, &argv);
   MPI_Comm_rank(MPI_COMM_WORLD, &MyProc);
   MPI_Comm_size(MPI_COMM_WORLD, &size);
   root = 0;
   if (MyProc == root)  /* El proceso raíz lee los límites del intervalo */
   { 
      printf("Give the left and right limits of the interval\n");
      scanf("%d %d", &left, &right);
      printf("Proc root reporting : the limits are : %d %d\n", left, right);
   }
   MPI_Bcast(&left, 1, MPI_INT, root, MPI_COMM_WORLD); /*Bcast limites a todos*/
   MPI_Bcast(&right, 1, MPI_INT, root, MPI_COMM_WORLD);
   if (((right - left + 1) % size) != 0)
      interval = (right - left + 1) / size + 1 ;  /*Fija límites locales de suma*/
   else
      interval = (right - left + 1) / size;
   mystart = left + MyProc*interval ;
   myend = mystart + interval ;
   /* establece los límites de los intervalos correctamente */
   if (myend > right) myend = right + 1 ;
   sum = root; /* Suma localmente en cada proceso MPI */
   if (mystart <= right)
      for (number = mystart; number < myend; number++) sum = sum + number ;
   /* Hace la reducción en el proceso raíz */
   MPI_Reduce(&sum, &GrandTotal, 1, MPI_INT, MPI_SUM, root, MPI_COMM_WORLD) ;
   MPI_Barrier(MPI_COMM_WORLD);
   /* El proceso raíz retorna los resultados */
   if(MyProc == root)
      printf("Proc root reporting : Grand total = %d \n", GrandTotal);
   MPI_Finalize();
}

int MPI_Allreduce (void *sendbuf, void *recvbuf, int count, 
            MPI_Datatype datatype, MPI_Op op, MPI_Comm comm)

   mpicc programa.c -o programa

   mpirun -np 4 programa

   mpirun -np 4 -machinefile nodes.txt programa

   node1
   node2

   int *p1;   /* puntero privado apuntando a datos locales */
   shared int *p2; /* puntero privado apuntando al espacio compartido */
   int *shared p3; /* puntero compartido apuntando a datos locales *;
   shared int *shared p4; /* puntero compartido apuntando al espacio compartido */

   shared int v1[N], v2[N], v1v2sum[N];

   void main()
   {
      int i;
      shared int *p1, *p2;
      p1=v1;
      p2=v2;
      upc_forall(i=0; i<N; i++; p1++; p2++; i)
      {
         v1v2sum[i]=*p1+*p2;
      }
   }

   real, dimension(n,n)[p,*]::a,b,c
   ...
   do k=1,n
      do q=1,p
              c(i,j)[myP,myQ]=c(i,j)[myP,myQ]+a(i,k)[myP,q]*b(k,j)[q,myQ]
      enddo
   enddo

public static void matMul(   double [2d] a,
               double [2d] b,
               double [2d] c)
{
   foreach (ij in c.domain())
   {
      double [1d] aRowi=a.slice(1, ij[1]);
      couble [1d] bColj=b.slice(2, ij[2]);
      foreach (k in aRowi.domain())
      {
         c[ij]+=aRowi[k]*bColj[k];
      }
   }
}

   int single stepCount=0;
   int single endCount=100;
   for (; stepCount<endCount; stepCount++)
   {
      Lee partículas remotas
      Calcula las fuerzas de la mía
      Ti.barrier();
      Escribe mis partículas utilizando nuevas fuerzas
      Ti.barrier();
   }