Ответ 1
ОБНОВЛЕНИЕ версия графического процессора
__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
}
}
Этот я только что тестировал на меньшие входные данные, потому что я тестирую свой ноутбук. Тем не менее, это сработало. Тем не менее, необходимо проводить дальнейшие испытания.
ОБНОВЛЕНИЕ Последовательная версия
Я просто сделал эту наивную версию, которая выполняет ваш алгоритм на 30 000 000 менее чем за 20 секунд (уже считая функцию для генерации данных).
В принципе, он сортирует ваш массив поплавков. Он будет перемещаться по отсортированному массиву, анализируя количество раз, которое значение последовательно появляется в массиве, а затем помещает это значение в словарь вместе с количеством раз, которое оно появляется.
Вы можете использовать отсортированную карту, а не неупорядоченный_мап, который я использовал.
Вот код:
#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;
}
Если у вас установлена установка библиотеки в вашем компьютере, вы должны использовать это:
#include <thrust/sort.h>
thrust::sort(data, data + size);
вместо этого
sort(data, data + size);
Конечно, это будет быстрее.
Оригинальная публикация
"Я работаю над статистическим приложением с большим массивом, содержащим 10-30 миллионов значений с плавающей запятой".
"Возможно ли (и имеет ли смысл) использовать графический процессор для ускорения таких вычислений?"
Да, это так. Месяц назад я поставил Molecular Dynamic полностью на GPU. Одно из ядер, которое вычисляет силу между парами частиц, получает 6 массивов, каждый из которых имеет 500 000 удвоений, в общей сложности 3 Миллиона удваивается (22 МБ).
Итак, вы планируете поставить 30 миллионов точек с плавающей точкой, это около 114 МБ глобальной памяти, поэтому это не проблема, даже мой ноутбук имеет 250 МБ.
В вашем случае может быть проблема с количеством вычислений? Основываясь на моем опыте с Molecular Dynamic (MD), я говорю "нет". Последовательная версия MD занимает около 25 часов, а в GPU - 45 минут. Вы сказали, что ваше приложение заняло пару часов, также в вашем примере кода оно выглядит мягче, чем Molecular Dynamic.
Здесь пример вычисления силы:
__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;
}
}
Простым примером кода в Cuda может быть сумма двух 2D-массивов:
В c:
for(int i = 0; i < N; i++)
c[i] = a[i] + b[i];
В 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];
}
В Cuda вы в основном брали каждую итерацию и делили по каждому потоку,
1) threadIdx.x + blockIdx.x*blockDim.x;
Каждый блок имеет идентификатор от 0 до N-1 (N максимальное число блоков), а каждый блок имеет число X потоков с идентификатором от 0 до X-1.
1) Дает вам итерацию, которую каждый поток будет вычислять на основе этого id и идентификатора блока, в котором находится поток, blockDim.x - это число потоков, которое имеет блок.
Итак, если у вас есть 2 блока, каждый из которых имеет 10 потоков и N = 40,
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
Глядя на ваш код, я сделал этот проект того, что может быть в 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();
}
Вы должны использовать atomicAdd, потому что разные потоки из разных блоков могут одновременно писать/читать noOfOccurrences, поэтому вам нужно неуверенно обойти исключение.
Это только один подход, который вы даже можете дать итерациям внешнего цикла для потоков вместо блоков.
Учебники
Серия журналов Dr Dobbs CUDA: Суперкомпьютер для масс Роб Фермера превосходна и охватывает почти все в четырнадцати партиях. Он также начинается довольно мягко и, следовательно, довольно дружелюбен для начинающих.
и другие:
- Разработка с CUDA - Введение
- Том I: Введение в программирование CUDA
- Начало работы с CUDA
- Список ресурсов CUDA
Взгляните на последний элемент, вы найдете много ссылок, чтобы узнать CUDA.
OpenCL: OpenCL Tutorials | MacResearch