c++ - tipos - ¿Puedo/debo ejecutar este código en una GPU?

tipos de gpu

Question

Estoy trabajando en una aplicación estadística que contiene aproximadamente 10 - 30 millones de valores de punto flotante en una matriz.

Varios métodos realizan cálculos diferentes, pero independientes, en la matriz en bucles anidados, por ejemplo:

Dictionary<float, int> noOfNumbers = new Dictionary<float, int>();

for (float x = 0f; x < 100f; x += 0.0001f) {
    int noOfOccurrences = 0;

    foreach (float y in largeFloatingPointArray) {
        if (x == y) {
            noOfOccurrences++;
        }
    }

    noOfNumbers.Add(x, noOfOccurrences);
}

La aplicación actual está escrita en C #, se ejecuta en una CPU Intel y necesita varias horas para completarse. No tengo conocimiento de los conceptos y API de programación de la GPU, por lo que mis preguntas son:

• ¿Es posible (y tiene sentido) utilizar una GPU para acelerar tales cálculos?
• En caso afirmativo: ¿Alguien conoce algún tutorial u obtiene un código de muestra (el lenguaje de programación no importa)?

Cualquier ayuda sería muy apreciada.

Answer 1

¿Es posible (y tiene sentido) utilizar una GPU para acelerar tales cálculos?

• Definitivamente SÍ , este tipo de algoritmo es, por lo general, el candidato ideal para el procesamiento masivo de paralelismo de datos, algo en lo que las GPU son tan buenas.

En caso afirmativo: ¿Alguien conoce algún tutorial u obtiene un código de muestra (el lenguaje de programación no importa)?

• Cuando desee utilizar GPGPU, tiene dos alternativas: CUDA o OpenCL .

  CUDA está maduro con muchas herramientas, pero las GPU de NVidia están centradas.

  OpenCL es un estándar que se ejecuta en las GPU NVidia y AMD, y en las CPU también. Entonces deberías realmente favorecerlo.

• Para el tutorial tiene una excelente serie sobre CodeProject por Rob Farber : http://www.codeproject.com/Articles/Rob-Farber#Articles

• Para su caso de uso específico, hay muchas muestras para construir histogramas con OpenCL (tenga en cuenta que muchos son histogramas de imágenes, pero los principios son los mismos).

• Al usar C #, puede usar enlaces como OpenCL.Net o Cloo .

• Si su matriz es demasiado grande para almacenarse en la memoria de la GPU, puede dividirla en particiones y volver a ejecutar su kernel OpenCL para cada parte fácilmente.

Answer 2

Además de la sugerencia del póster anterior, use TPL (biblioteca paralela de tareas) cuando sea apropiado para ejecutar en paralelo en múltiples núcleos.

El ejemplo anterior podría usar Parallel.Foreach y ConcurrentDictionary, pero un mapa más complejo: reduzca la configuración donde la matriz se divide en fragmentos, cada uno de los cuales genera un diccionario que luego se reduciría a un solo diccionario le ofrecería mejores resultados.

No sé si todos sus cálculos se correlacionan correctamente con las capacidades de la GPU, pero tendrá que usar un algoritmo de reducción de mapa de todos modos para asignar los cálculos a los núcleos de la GPU y luego reducir los resultados parciales a un solo resultado, para que pueda también podría hacerlo en la CPU antes de pasar a una plataforma menos familiar.

Answer 3

No estoy seguro de si usar GPU sería una buena combinación dado que los valores ''largerFloatingPointArray'' necesitan recuperarse de la memoria. Tengo entendido que las GPU son más adecuadas para cálculos autocontenidos.

Creo que convertir esta aplicación de proceso único en una aplicación distribuida que se ejecuta en muchos sistemas y modificar el algoritmo debería acelerar considerablemente, dependiendo de cuántos sistemas estén disponibles.

Puedes usar el enfoque clásico de ''divide y vencerás''. El enfoque general que tomaría es el siguiente.

Use un sistema para preprocesar ''largeFloatingPointArray'' en una tabla hash o una base de datos. Esto se haría en una sola pasada. Utilizaría el valor del punto flotante como la clave y el número de apariciones en la matriz como el valor. El peor escenario es que cada valor solo se produce una vez, pero eso es poco probable. Si largeFloatingPointArray cambia constantemente cada vez que se ejecuta la aplicación, entonces la tabla hash en memoria tiene sentido. Si es estático, la tabla podría guardarse en una base de datos de valores clave como Berkeley DB. Llamemos a esto un sistema de "búsqueda".

En otro sistema, vamos a llamarlo ''principal'', crear trozos de trabajo y ''dispersar'' los elementos de trabajo en N sistemas, y ''reunir'' los resultados a medida que estén disponibles. Por ejemplo, un elemento de trabajo podría ser tan simple como dos números que indiquen el rango en el que debería funcionar un sistema. Cuando un sistema completa el trabajo, envía una serie de incidencias y está listo para trabajar en otro trozo de trabajo.

El rendimiento se mejora porque no seguimos iterando sobre largeFloatingPointArray. Si el sistema de búsqueda se convierte en un cuello de botella, se podría replicar en tantos sistemas como sea necesario.

Con un número suficientemente grande de sistemas trabajando en paralelo, debería ser posible reducir el tiempo de procesamiento a minutos.

Estoy trabajando en un compilador para programación paralela en C para sistemas basados ​​en muchos núcleos, a menudo conocidos como microservidores, que se construirán usando múltiples módulos ''system-on-a-chip'' dentro de un sistema. Los proveedores de módulos ARM incluyen Calxeda, AMD, AMCC, etc. Intel probablemente también tenga una oferta similar.

Tengo una versión del compilador funcionando, que podría usarse para esa aplicación. El compilador, basado en prototipos de función C, genera un código de red C que implementa el código de comunicación entre procesos (IPC) en todos los sistemas. Uno de los mecanismos de IPC disponibles es socket / tcp / ip.

Si necesita ayuda para implementar una solución distribuida, me gustaría discutirla con usted.

Agregado el 16 de noviembre de 2012.

Pensé un poco más sobre el algoritmo y creo que esto debería hacerlo en una sola pasada. Está escrito en C y debería ser muy rápido comparado con lo que tienes.

/*
 * Convert the X range from 0f to 100f in steps of 0.0001f
 * into a range of integers 0 to 1 + (100 * 10000) to use as an
 * index into an array.
 */

#define X_MAX           (1 + (100 * 10000))

/*
 * Number of floats in largeFloatingPointArray needs to be defined
 * below to be whatever your value is.
 */

#define LARGE_ARRAY_MAX (1000)

main()
{
    int j, y, *noOfOccurances;
    float *largeFloatingPointArray;

    /*
     * Allocate memory for largeFloatingPointArray and populate it.
     */

    largeFloatingPointArray = (float *)malloc(LARGE_ARRAY_MAX * sizeof(float));    
    if (largeFloatingPointArray == 0) {
        printf("out of memory/n");
        exit(1);
    }

    /*
     * Allocate memory to hold noOfOccurances. The index/10000 is the
     * the floating point number.  The contents is the count.
     *
     * E.g. noOfOccurances[12345] = 20, means 1.2345f occurs 20 times
     * in largeFloatingPointArray.
     */

    noOfOccurances = (int *)calloc(X_MAX, sizeof(int));
    if (noOfOccurances == 0) {  
        printf("out of memory/n");
        exit(1);
    }

    for (j = 0; j < LARGE_ARRAY_MAX; j++) {
        y = (int)(largeFloatingPointArray[j] * 10000);
        if (y >= 0 && y <= X_MAX) {
            noOfOccurances[y]++;
        }   
    }
}

Answer 4

No sé mucho sobre el procesamiento en paralelo o GPGPU, pero para este ejemplo específico, podría ahorrar mucho tiempo haciendo una sola pasada sobre la matriz de entrada en lugar de recorrerla un millón de veces. Con grandes conjuntos de datos, normalmente querrá hacer cosas en un solo pase si es posible. Incluso si realiza múltiples cálculos independientes, si supera el mismo conjunto de datos, es posible que obtenga una mayor velocidad haciendo todos los pasos en el mismo pase, ya que obtendrá una mejor localidad de referencia de esa manera. Pero puede no valer la pena por la mayor complejidad en su código.

Además, realmente no desea agregar una cantidad pequeña a un número de coma flotante de manera repetitiva, el error de redondeo se sumará y no obtendrá lo que deseaba. He agregado una declaración if a mi ejemplo siguiente para verificar si las entradas coinciden con tu patrón de iteración, pero omítalo si realmente no lo necesitas.

No conozco ningún C #, pero una implementación de una sola vez de su muestra se vería así:

Dictionary<float, int> noOfNumbers = new Dictionary<float, int>();

foreach (float x in largeFloatingPointArray)
{
    if (math.Truncate(x/0.0001f)*0.0001f == x)
    {
        if (noOfNumbers.ContainsKey(x))
            noOfNumbers.Add(x, noOfNumbers[x]+1);
        else
            noOfNumbers.Add(x, 1);
    }
}

Espero que esto ayude.

Answer 5

ACTUALIZAR la versión de la GPU

__global__ void hash (float *largeFloatingPointArray,int largeFloatingPointArraySize, int *dictionary, int size, int num_blocks)
{
    int x = (threadIdx.x + blockIdx.x * blockDim.x); // Each thread of each block will
    float y;                                         // compute one (or more) floats
    int noOfOccurrences = 0;
    int a;

    while( x < size )            // While there is work to do each thread will:
    {
        dictionary[x] = 0;       // Initialize the position in each it will work
        noOfOccurrences = 0;    

        for(int j = 0 ;j < largeFloatingPointArraySize; j ++) // Search for floats
        {                                                     // that are equal 
                                                             // to it assign float
           y = largeFloatingPointArray[j];  // Take a candidate from the floats array 
           y *= 10000;                      // e.g if y = 0.0001f;
           a = y + 0.5;                     // a = 1 + 0.5 = 1;
           if (a == x) noOfOccurrences++;    
        }                                      

        dictionary[x] += noOfOccurrences; // Update in the dictionary 
                                          // the number of times that the float appears 

    x += blockDim.x * gridDim.x;  // Update the position here the thread will work
    }
}

Este último lo probé para entradas más pequeñas, porque estoy probando mi computadora portátil. Sin embargo, funcionó. Sin embargo, es necesario hacer más testículos.

ACTUALIZAR la versión secuencial

Acabo de hacer esta versión ingenua que realiza su algoritmo por 30,000,000 en menos de 20 segundos (ya cuenta la función para generar datos).

Básicamente, clasifica tu conjunto de flotadores. Recorrerá la matriz ordenada, analizando el número de veces que un valor aparece consecutivamente en la matriz y luego coloca este valor en un diccionario junto con el número de veces que aparece.

Puede usar el mapa ordenado, en lugar del mapa desordenado que utilicé.

Aquí está el código:

#include <stdio.h>
#include <stdlib.h>
#include "cuda.h"
#include <algorithm>
#include <string>
#include <iostream>
#include <tr1/unordered_map>


typedef std::tr1::unordered_map<float, int> Mymap;


void generator(float *data, long int size)
{
    float LO = 0.0;
    float HI = 100.0;

    for(long int i = 0; i < size; i++)
        data[i] = LO + (float)rand()/((float)RAND_MAX/(HI-LO));
}

void print_array(float *data, long int size)
{

    for(long int i = 2; i < size; i++)
        printf("%f/n",data[i]);

}

std::tr1::unordered_map<float, int> fill_dict(float *data, int size)
{
    float previous = data[0];
    int count = 1;
    std::tr1::unordered_map<float, int> dict;

    for(long int i = 1; i < size; i++)
    {
        if(previous == data[i])
            count++;
        else
        {
          dict.insert(Mymap::value_type(previous,count));
          previous = data[i];
          count = 1;         
        }

    }
    dict.insert(Mymap::value_type(previous,count)); // add the last member
    return dict;

}

void printMAP(std::tr1::unordered_map<float, int> dict)
{
   for(std::tr1::unordered_map<float, int>::iterator i = dict.begin(); i != dict.end(); i++)
  {
     std::cout << "key(string): " << i->first << ", value(int): " << i->second << std::endl;
   }
}


int main(int argc, char** argv)
{
  int size = 1000000; 
  if(argc > 1) size = atoi(argv[1]);
  printf("Size = %d",size);

  float data[size];
  using namespace __gnu_cxx;

  std::tr1::unordered_map<float, int> dict;

  generator(data,size);

  sort(data, data + size);
  dict = fill_dict(data,size);

  return 0;
}

Si tiene instalado el empuje de la biblioteca en su máquina, debe usar esto:

#include <thrust/sort.h>
thrust::sort(data, data + size);

en lugar de esto

sort(data, data + size);

Por supuesto, será más rápido.

Publicación original

"Estoy trabajando en una aplicación estadística que tiene una gran matriz que contiene entre 10 y 30 millones de valores de coma flotante".

"¿Es posible (y tiene sentido) utilizar una GPU para acelerar tales cálculos?"

Sí lo es. Hace un mes puse una simulación Molecular Dynamic completamente en la GPU. Uno de los núcleos, que calcula la fuerza entre pares de partículas, recibe 6 arreglos cada uno con 500,000 dobles, un total de 3 Millones de dobles (22 MB).

Así que está planeando poner 30 millones de puntos flotantes, esto es alrededor de 114 MB de memoria global, así que esto no es un problema, incluso mi computadora portátil tiene 250 MB.

El número de cálculos puede ser un problema en su caso? Basado en mi experiencia con la Molecular Dynamic (MD) digo que no. La versión secuencial de MD tarda aproximadamente 25 horas en completarse mientras que en la GPU tomó 45 minutos. Dijiste que tu aplicación tomó un par de horas, también se basa en el ejemplo de tu código, se ve más suave que el Molecular Dynamic.

Aquí está el ejemplo de cálculo de fuerza:

__global__ void add(double *fx, double *fy, double *fz,
                    double *x, double *y, double *z,...){

     int pos = (threadIdx.x + blockIdx.x * blockDim.x); 

     ...

     while(pos < particles)
     {

      for (i = 0; i < particles; i++)
      {
              if(//inside of the same radius)
                {
                 // calculate force
                } 
       }
     pos += blockDim.x * gridDim.x;  
     }        
  }

Un ejemplo simple de un código en Cuda podría ser la suma de dos matrices 2D:

Cª:

for(int i = 0; i < N; i++)
    c[i] = a[i] + b[i]; 

En Cuda:

__global__ add(int *c, int *a, int*b, int N)
{
  int pos = (threadIdx.x + blockIdx.x)
  for(; i < N; pos +=blockDim.x)
      c[pos] = a[pos] + b[pos];
}

En Cuda, básicamente tomaste cada una para la iteración y dividiste por cada hilo,

1) threadIdx.x + blockIdx.x*blockDim.x;

Cada bloque tiene una Id de 0 a N-1 (N es el número máximo de bloques) y cada bloque tiene un número X de hilos con una identificación de 0 a X-1.

1) Te da la iteración de iteración que cada subproceso calculará en función de su id. Y el id. De bloque donde está el subproceso, blockDim.x es el número de subproceso que tiene un bloque.

Entonces, si tienes 2 bloques cada uno con 10 hilos y un N = 40, el:

Thread 0 Block 0 will execute pos 0
Thread 1 Block 0 will execute pos 1
...
Thread 9 Block 0 will execute pos 9
Thread 0 Block 1 will execute pos 10
....
Thread 9 Block 1 will execute pos 19
Thread 0 Block 0 will execute pos 20
...
Thread 0 Block 1 will execute pos 30
Thread 9 Block 1 will execute pos 39

Mirando su código hice este borrador de lo que podría ser en cuda:

__global__ hash (float *largeFloatingPointArray, int *dictionary)
    // You can turn the dictionary in one array of int
    // here each position will represent the float
    // Since  x = 0f; x < 100f; x += 0.0001f
    // you can associate each x to different position
    // in the dictionary:

    // pos 0 have the same meaning as 0f;
    // pos 1 means float 0.0001f
    // pos 2 means float 0.0002f ect.
    // Then you use the int of each position 
    // to count how many times that "float" had appeared 


   int x = blockIdx.x;  // Each block will take a different x to work
    float y;

while( x < 1000000) // x < 100f (for incremental step of 0.0001f)
{
    int noOfOccurrences = 0;
    float z = converting_int_to_float(x); // This function will convert the x to the
                                          // float like you use (x / 0.0001)

    // each thread of each block
    // will takes the y from the array of largeFloatingPointArray

    for(j = threadIdx.x; j < largeFloatingPointArraySize; j += blockDim.x)
    {
        y = largeFloatingPointArray[j];
        if (z == y)
        {
            noOfOccurrences++;
        }
    }
    if(threadIdx.x == 0) // Thread master will update the values
      atomicAdd(&dictionary[x], noOfOccurrences);
    __syncthreads();
}

Debe utilizar atomicAdd porque diferentes hilos de diferentes bloques pueden escribir / leer noOfOccurrences al mismo tiempo, por lo que debe asegurarse de la exclusión mutua.

Este es solo un enfoque, incluso puede dar las iteraciones del bucle externo a los hilos en lugar de los bloques.

Tutoriales

La serie del Dr. Dobbs Journal CUDA: la supercomputación para las masas por Rob Farmer es excelente y abarca casi todo en sus catorce cuotas. También comienza con bastante suavidad y, por lo tanto, es bastante amigable para principiantes.

y anothers:

• Desarrollando con CUDA - Introducción
• Volumen I: Introducción a la Programación CUDA
• Comenzando con CUDA
• Lista de recursos de CUDA

Echa un vistazo al último artículo, encontrarás muchos enlaces para aprender CUDA.

OpenCL: Tutoriales de OpenCL | MacResearch