Может/должен ли я запускать этот код на графическом процессоре?

Ответ 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

Ответ 2

Я ничего не знаю о параллельной обработке или GPGPU, но для этого конкретного примера вы могли бы сэкономить много времени, сделав один проход над входным массивом, а не зацикливая на него миллион раз. С большими наборами данных вы, как правило, захотите сделать что-то за один проход, если это возможно. Даже если вы выполняете несколько независимых вычислений, если это над одним и тем же набором данных, вы можете получить более высокую скорость, делая все это в одном и том же проходе, так как вы получите лучшую локальность ссылок. Но это может не стоить того, чтобы повысить сложность вашего кода.

Кроме того, вы действительно не хотите добавлять небольшую сумму в число с плавающей запятой, повторяющееся так, ошибка округления будет складываться, и вы не получите того, что вы намеревались. Я добавил оператор if в свой пример ниже, чтобы проверить, соответствуют ли входы вашему шаблону итерации, но опустите его, если вам это действительно не нужно.

Я не знаю ни одного С#, но реализация одного прохода вашего образца будет выглядеть примерно так:

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);
    }
}

Надеюсь, что это поможет.

Ответ 3

Возможно ли (и имеет смысл) использовать графический процессор для ускорения такие вычисления?

• Определенно ДА, этот тип алгоритма, как правило, является идеальным кандидатом для массивной обработки данных parallelism, то, что GPU настолько хороши.

Если да: кто-нибудь знает какой-либо учебник или получает какой-либо пример кода (язык программирования не имеет значения)?

• Если вы хотите перейти по пути GPGPU, у вас есть две альтернативы: CUDA или OpenCL.

  CUDA является зрелым с большим количеством инструментов, но графический процессор NVidia ориентирован.

  OpenCL - это стандартная работа на графических процессорах NVidia и AMD, а также на процессорах. Так что вы действительно должны это одобрить.

• Для учебника у вас есть отличная серия по CodeProject Роба Фарбера: http://www.codeproject.com/Articles/Rob-Farber#Articles

• В вашем конкретном случае использования существует множество образцов для гистограмм, построенных с помощью OpenCL (обратите внимание, что многие из них являются гистограммами изображений, но принципы одинаковы).

• Как вы используете С#, вы можете использовать привязки типа OpenCL.Net или Cloo.

• Если ваш массив слишком велик для хранения в памяти графического процессора, вы можете заблокировать его и повторно запустить ядро ​​OpenCL для каждой части.

Ответ 4

В дополнение к предложению, приведенному выше, используйте TPL (параллельную библиотеку задач), когда это необходимо, для параллельной работы на нескольких ядрах.

В приведенном выше примере можно использовать Parallel.Foreach и ConcurrentDictionary, но более сложная установка с уменьшением карты, где массив разбивается на куски, каждый из которых генерирует словарь, который затем сводится к одному словарю, даст вам лучшие результаты.

Я не знаю, правильно ли все ваши вычисления правильно отображают возможности графического процессора, но вам все равно придется использовать алгоритм сокращения карты, чтобы сопоставить вычисления с ядрами GPU, а затем уменьшить частичные результаты до одного результата, поэтому вы можете также сделать это на процессоре, прежде чем перейти к менее знакомой платформе.

Ответ 5

Я не уверен, будет ли использование графических процессоров хорошим совпадением, учитывая, что Значения "largeFloatingPointArray" необходимо извлекать из памяти. Я понимаю, что графические процессоры лучше подходят для самостоятельных вычислений.

Я думаю, что превратить это одно приложение процесса в распределенное приложение, работающее на многих системах, и настройка алгоритма должна значительно ускорить процесс, в зависимости от того, сколько систем доступно.

Вы можете использовать классический подход "разделяй и властвуй". Общий подход, который я хотел бы сделать, заключается в следующем.

Используйте одну систему для предварительной обработки 'largeFloatingPointArray' в хеш-таблице или базе данных. Это было бы сделано за один проход. Он будет использовать значение с плавающей запятой в качестве ключа и количество вхождений в массиве в качестве значения. Худший сценарий - это то, что каждое значение возникает только один раз, но это маловероятно. Если largeFloatingPointArray постоянно меняется при каждом запуске приложения, то в хэш-таблице в памяти имеет смысл. Если он статичен, таблица может быть сохранена в базе данных с ключом, такой как Berkeley DB. Позвольте называть это системой поиска.

В другой системе позвольте ей "main", создать куски работы и "разбросать" рабочие элементы по N системам и "собрать" результаты по мере их появления. Например, рабочий элемент может быть таким же простым, как два числа, указывающие диапазон, над которым система должна работать. Когда система завершает работу, она отправляет обратно массив вхождений и готова к работе над другим фрагментом работы.

Производительность улучшена, потому что мы не выполняем итерацию над largeFloatingPointArray. Если система поиска становится узким местом, тогда она может быть реплицирована на столько систем, сколько необходимо.

При достаточно большом количестве систем, работающих параллельно, должно быть возможно сократить время обработки до минут.

Я работаю над компилятором для параллельного программирования в C, предназначенным для многоядерных систем, часто называемых микросервистами, которые будут/или будут построены с использованием нескольких модулей "система на кристалле" в системе, Поставщики модулей ARM включают Calxeda, AMD, AMCC и т.д. У Intel, вероятно, также будет аналогичное предложение.

У меня есть версия компилятора, которая может быть использована для такого приложения. Компилятор, основанный на прототипах функций C, генерирует сетевой код C, который реализует межпроцессный коммуникационный код (IPC) в разных системах. Одним из доступных механизмов IPC является socket/tcp/ip.

Если вам нужна помощь в реализации распределенного решения, я буду рад обсудить его с вами.

Добавлено 16 ноября 2012 г.

Я подумал немного больше об алгоритме, и я думаю, что это должно сделать это за один проход. Он написан на C, и он должен быть очень быстрым по сравнению с тем, что у вас есть.

/*
 * 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]++;
        }   
    }
}