可以/我应该在GPU上运行此代码吗？

Question

我正在研究一个统计应用程序,它在一个数组中包含大约10-30万个浮点值.

有几种方法在嵌套循环中对数组执行不同但独立的计算,例如:

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

当前的应用程序是用C#编写的,在Intel CPU上运行,需要几个小时才能完成.我不了解GPU编程概念和API,所以我的问题是:

• 是否有可能(并且有意义)利用GPU来加速这样的计算？
• 如果是:有没有人知道任何教程或获得任何示例代码(编程语言无关紧要)？

任何帮助将受到高度赞赏.

Answer 1

更新 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
    }
}

这个我刚刚测试了较小的输入,因为我正在测试我的笔记本电脑.尽管如此,它确实有效.但是,有必要做更进一步的睾丸.

更新顺序版本

我只是做了这个天真的版本,在不到20秒的时间内执行你的算法30,000,000(已经计算生成数据的函数).

基本上,它会对您的浮动数组进行排序.它将遍历排序的数组,分析值连续出现在数组中的次数,然后将此值与其出现的次数一起放入字典中.

您可以使用有序映射,而不是我使用的unordered_map.

下面是代码:

#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百万个浮点值的大数组".

"利用GPU来加速这样的计算是否可能(并且有意义)？"

是的.一个月前,我将分子动态模拟完全放在GPU上.计算粒子对之间的力的核心之一是每个接收6个阵列,500,000个双打,总计3百万双(22 MB).

因此,您计划投入30百万浮点数,这大约是114 MB的全局内存,所以这不是问题,即使我的笔记本电脑有250MB.

在您的情况下,计算的数量可能是一个问题？根据我对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]; 

在库达:

__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的Id(N是块的最大数量),并且每个块具有X个线程,其ID从0到X-1.

1)给出每个线程将基于它的id和线程所在的块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 Journal系列CUDA: Rob Farmer为大众提供的超级计算非常出色,涵盖了十四期中的所有内容.它也开始相当温和,因此相当初学者友好.

和其他人:

• 用CUDA开发 - 介绍
• 第一卷:CUDA编程简介
• 开始使用CUDA
• CUDA资源清单

看看最后一项,你会发现许多学习CUDA的链接.

OpenCL:OpenCL教程| MacResearch

Answer 2

我对并行处理或GPGPU一无所知,但对于这个具体的例子,你可以通过对输入数组进行一次传递而不是循环遍历它一百万次来节省大量时间.对于大型数据集,如果可能,您通常希望一次性完成任务.即使您正在进行多次独立计算,如果它在同一个数据集上,您可能会在同一个传递中更快地完成所有这些计算,因为您将获得更好的参考位置.但是,对于代码中复杂性的增加,它可能不值得.

另外,你真的不想像浮动数字那样重复添加少量数据,舍入错误会加起来,你不会得到你想要的.我在下面的示例中添加了一个if语句,以检查输入是否与您的迭代模式匹配,但如果您实际上并不需要,则省略它.

我不知道任何C#,但您的示例的单个传递实现看起来像这样:

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

希望这可以帮助.

Answer 3

是否有可能(并且有意义)利用GPU来加速这样的计算？

• 绝对是的,这种算法通常是大规模数据并行处理的理想选择,GPU非常擅长.

如果是:有没有人知道任何教程或获得任何示例代码(编程语言无关紧要)？

• 当您想要采用GPGPU方式时,您有两种选择:CUDA或OpenCL.

  CUDA已经成熟,有很多工具,但是以NVidia GPU为中心.

  OpenCL是在NVidia和AMD GPU以及CPU上运行的标准.所以你应该真的喜欢它.

• 在教程中,你有一篇关于Rob Farber的 CodeProject的精彩系列:http://www.codeproject.com/Articles/Rob-Farber#Articles

• 对于您的特定用例,有许多直方图与OpenCL建立的样本(请注意,许多是图像直方图,但原理是相同的).

• 当您使用C#时,您可以使用OpenCL.Net或Cloo等绑定.

• 如果您的阵列太大而无法存储在GPU内存中,您可以对其进行分区分区,并轻松地为每个部分重新运行OpenCL内核.

Answer 4

除了上述海报的建议之外,在适当的情况下使用TPL(任务并行库)在多个核上并行运行.

上面的示例可以使用Parallel.Foreach和ConcurrentDictionary,但是更复杂的map-reduce设置,其中数组被拆分为块,每个块生成一个字典,然后将其缩减为单个字典将为您提供更好的结果.

我不知道您的所有计算是否都能正确映射到GPU功能,但您无论如何都必须使用map-reduce算法将计算映射到GPU核心,然后将部分结果减少为单个结果,因此您在进入一个不太熟悉的平台之前,不妨在CPU上做到这一点.

Answer 5

我不确定使用GPU是否是一个很好的匹配,因为需要从内存中检索'greaterFloatingPointArray'值.我的理解是GPU更适合自包含计算.

我认为将这个单一流程应用程序转换为在许多系统上运行的分布式应用程序并调整算法应该会大大加快速度,具体取决于可用的系统数量.

您可以使用经典的"分而治之"方法.我将采取的一般方法如下.

使用一个系统将'largeFloatingPointArray'预处理为哈希表或数据库.这将在一次通过中完成.它将使用浮点值作为键,并将数组中出现的次数作为值.最糟糕的情况是每个值只出现一次,但这不太可能.如果每次运行应用程序时,largeFloatingPointArray都会不断更改,那么内存中的哈希表就有意义了.如果它是静态的,则表可以保存在键值数据库中,例如Berkeley DB.我们称之为"查找"系统.

在另一个系统中,让我们将其称为"主要",创建工作块并在N个系统中"分散"工作项,并在结果可用时"收集"结果.例如,工作项可以像两个数字一样简单,表示系统应该处理的范围.当系统完成工作时,它会发回一系列事件,并准备好处理另一部分工作.

性能得到改善,因为我们不会继续迭代largeFloatingPointArray.如果查找系统成为瓶颈,那么它可以根据需要在尽可能多的系统上进行复制.

有足够多的系统并行工作,应该可以将处理时间缩短到几分钟.

我正在开发一种用于C语言并行编程的编译器,用于基于多核的系统(通常称为微服务器),这些系统是/或将使用系统内的多个"片上系统"模块构建.ARM模块供应商包括Calxeda,AMD,AMCC等.英特尔可能也会提供类似的产品.

我有一个版本的编译器工作,可用于这样的应用程序.该编译器基于C函数原型生成C网络代码,该代码跨系统实现进程间通信代码(IPC).可用的IPC机制之一是socket/tcp/ip.

如果您在实施分布式解决方案时需要帮助,我很乐意与您讨论.

添加2012年11月16日.

我想到了更多关于算法的内容,我认为这应该在一次通过中完成.它是用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]++;
        }   
    }
}