43

我正在开发一个统计应用程序,该应用程序在一个数组中包含大约 10-3000 万个浮点值。

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

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 来加速此类计算?
  • 如果是:有谁知道任何教程或有任何示例代码(编程语言无关紧要)?
4

5 回答 5

88

更新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 个元素的数组执行算法。

这个幼稚的版本首先对您的浮点数组进行排序。之后,将遍历已排序的数组并检查给定value值在数组中出现的次数,然后将该值连同它出现的次数一起放入字典中。

您可以使用sorted地图,而不是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是全局内存,那不是问题。

在您的情况下,计算次数会成为问题吗?根据我在分子动力学 (MD) 方面的经验,我会说不。顺序 MD 版本大约需要25几个小时才能完成,而 GPU 版本需要45几分钟。您说您的应用程序需要几个小时,同样基于您的代码示例,它看起来比 MD更软。

这是力计算示例:

__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 中代码的一个简单示例可以是两个二维数组的总和:

在 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;

每个块都有一个IDfrom 0to N-1(N 块的最大数量),每个块都有'X'多个线程,ID0to X-1

  1. 为您提供每个线程将根据其和线程所在块计算的for循环迭代;blockDim.x 是块拥有的线程数。IDID

因此,如果您有 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同时写入/读取,因此您必须确保互斥

这只是一种方法;您甚至可以将外部循环的迭代分配给线程而不是块。

教程

Dobbs 博士期刊系列CUDA:Rob Farmer 的大众超级计算非常出色,在其十四期中几乎涵盖了所有内容。它也开始相当温和,因此对初学者相当友好。

和其他人:

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

OpenCL:OpenCL 教程 | 麦克研究

于 2012-11-15T07:24:10.250 回答
11

我对并行处理或 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);
    }
}

希望这可以帮助。

于 2012-11-09T03:37:54.943 回答
9

是否有可能(并且有意义)利用 GPU 来加速此类计算?

  • 绝对的,这种算法通常是海量数据并行处理的理想候选者,GPU 非常擅长的事情。

如果是:有谁知道任何教程或有任何示例代码(编程语言无关紧要)?

  • 当您想采用 GPGPU 方式时,您有两种选择:CUDAOpenCL

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

    OpenCL 是在 NVidia 和 AMD GPU 以及 CPU 上运行的标准。所以你真的应该赞成它。

  • 对于教程,您有Rob Farber的 CodeProject 优秀系列:http : //www.codeproject.com/Articles/Rob-Farber#Articles

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

  • 当您使用 C# 时,您可以使用OpenCL.NetCloo等绑定。

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

于 2012-11-13T09:56:06.693 回答
6

除了上述海报的建议外,在适当的时候使用 TPL(任务并行库)在多个内核上并行运行。

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

我不知道您的所有计算是否正确映射到 GPU 功能,但无论如何您都必须使用 map-reduce 算法将计算映射到 GPU 核心,然后将部分结果减少为单个结果,所以您在转移到不太熟悉的平台之前,不妨在 CPU 上执行此操作。

于 2012-11-09T03:49:05.993 回答
6

鉴于需要从内存中检索“largerFloatingPointArray”值,我不确定使用 GPU 是否是一个很好的匹配。我的理解是 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]++;
        }   
    }
}
于 2012-11-12T05:13:37.597 回答