Mise à JOUR Version 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
}
}
Celui-ci, je viens de tester pour les petites entrées, parce que je suis en train de tester, j'ai mon ordinateur portable. Néanmoins, il a fait un travail. Cependant, il est nécessaire de faire favorise les testicules.
Mise à JOUR de la Version Séquentielle
J'ai juste fait cette version naïve que d'effectuer votre algorithme pour de 30 000 000 en moins de 20 secondes (en comptant déjà de fonction pour générer des données).
En gros, c'est un peu votre tableau de float. Il sera du voyage sur le tableau trié, en analysant le nombre de fois qu'une valeur consécutivement apparaît dans le tableau et ensuite mettre cette valeur dans un dictionnaire, avec le nombre de fois où il apparaît.
Vous pouvez utiliser triés carte, au lieu de la unordered_map que j'ai utilisé.
Voici le code:
#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 vous avez de la bibliothèque de poussée installés dans votre machine, vous devez utiliser ceci:
#include <thrust/sort.h>
thrust::sort(data, data + size);
au lieu de cela
sort(data, data + size);
Pour sûr, ça sera plus rapide.
Post Original
"Je travaille sur une application de la statistique, qui dispose d'un large éventail containin 10 - 30 millions de valeurs à virgule flottante".
"Est-il possible (et du sens) d'utiliser un GPU pour accélérer les calculs?"
Oui, il est. Il y a un mois, j'ai mis une simulation Dynamique Moléculaire entièrement sur le GPU. L'un des noyaux, qui calcule la force entre les paires de particules, de recevoir 6 de la matrice de chacune de 500 000 chambres doubles, un total de 3 Millions doubles (22 MO).
Si vous êtes rabotage à mettre 30 Millions de flotteur de points est d'environ 114 MO de Mémoire globale, donc ce n'est pas un problème, même mon portable ont 250MB.
Le nombre de calcul peut être un problème dans votre cas? Basé sur mon expérience avec la Dynamique Moléculaire (MD) je dis non. Le séquentiel MD version prend environ 25 heures pour terminer tandis que dans le GPU a duré 45 Minutes. Vous avez dit que votre demande a fallu quelques heures, également basé dans le code de l'exemple, il semble plus doux que la Dynamique Moléculaire.
Voici la force de l'exemple de calcul:
__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 exemple simple de code Cuda peut être la somme de deux tableaux 2D:
En c:
for(int i = 0; i < N; i++)
c[i] = a[i] + b[i];
Dans 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];
}
Dans Cuda vous a pris pour chaque itération et de diviser par chaque thread,
1) threadIdx.x + blockIdx.x*blockDim.x;
Chaque bloc ont un Id de 0 à N-1 (N le nombre maximum de blocs) et chaque bloc ont un X nombre de threads avec un id de 0 à X-1.
1) vous Donne la pour l'itération que chaque thread de calcul basé sur l'id et le bloc id où le fil est dans, la blockDim.x est le nombre de threads d'un bloc ont.
Donc si vous avez 2 pâtés de maisons, chacune avec 10 threads et N = 40, l:
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
À la recherche de votre code, j'ai fait cette ébauche de ce que pourrait être dans 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();
}
Vous devez utiliser atomicAdd parce que les différents threads de différents blocs peut écrire/lire noOfOccurrences en même temps, si vous avez des doutes exclusion mutuelle.
Ce n'est qu'une approche, vous pouvez même donner les itérations de la boucle externe pour les threads au lieu de blocs.
Tutoriels
Le Dr Dobbs Journal de la série CUDA: calcul intensif pour les masses par Rob Agriculteur est excellent et couvre à peu près tout dans ses quatorze ans de versements. Elle commence plutôt doucement et est donc assez débutant-friendly.
et d'autres:
Prendre un coup d'oeil sur la dernière rubrique, vous trouverez de nombreux lien pour en savoir CUDA.
OpenCL: OpenCL Tutoriels | MacResearch
