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

我正在处理一个统计应用程序,其中包含大约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来加速这样的计算是否可能(并且是否有意义)?
  • 如果是的话:有谁知道任何教程或得到任何示例代码(编程语言无所谓)?

任何帮助将不胜感激。

更新 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; // eg 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 } } 

这个我刚刚testing了较小的input,因为我正在testing我的笔记本电脑。 不过,它确实有效。 但是,还有必要做进一步的睾丸检查。

更新顺序版本

我只是做了这个天真的版本,执行您的algorithm在不到20秒的30,000,000(已经计算function生成数据)。

基本上,它sorting你的浮点数组。 它将遍历已sorting的数组,分析数组中连续出现的值的次数,然后将该值与其出现的次数一起放入字典中。

你可以使用sorting的地图,而不是我使用的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); 

肯定会更快。

原始post

“我正在做一个统计应用程序,这个应用程序有一个包含10到30百万个浮点值的大数组”。

“是否有可能(而且是否有意义)利用GPU来加速这样的计算?”

是的。 一个月前,我完全在GPU上进行了分子动力学模拟。 内核之一计算粒子对之间的力,每个arrays接受6个arrays,每个arrays有500,000个双打,共有3百万双打(22 MB)。

所以你打算把30百万浮点数,这是大约114 MB的全球内存,所以这不是一个问题,即使我的笔记本电脑有250MB。

计算的数量可能是一个问题在你的情况? 基于我对分子动力学(MD)的经验,我说不。 顺序MD版本需要大约25小时才能完成,而在GPU中需要45分钟。 你说你的应用程序花了几个小时,也是基于你的代码示例,它看起来比分子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中一个简单的代码示例可以是两个二维数组的总和:

在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的Id(N是块的最大数量),每个块都有X个线程,ID从0到X-1。

1)为每个线程根据id和线程所在的块ID计算迭代,blockDim.x是一个块所具有的线程数。

所以如果你有两个块,每个块有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进行开发 – 介绍
  • 第一卷:CUDA编程简介
  • 开始使用CUDA
  • CUDA资源列表

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

OpenCL: OpenCL教程| MacResearch

对于并行处理或GPGPU,我并不了解太多,但是对于这个特定的例子,通过在input数组上进行一次遍历而不是循环遍历数百万次就可以节省大量的时间。 对于大数据集,如果可能的话,您通常会希望一次执行一些操作。 即使你正在做多个独立的计算,如果它是在同一个数据集上,你可能会得到更好的速度,在同一遍中完成它们,因为这样可以获得更好的参考局部性。 但是,对于代码中增加的复杂性,这可能不值得。

另外,你真的不想重复添加一个浮点数的小数,舍入误差会加起来,你不会得到你想要的。 我在下面的示例中添加了一个if语句来检查input是否符合迭代模式,但如果实际上并不需要,则省略它。

我不知道任何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); } } 

希望这可以帮助。

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

  • 肯定是的 ,这种algorithm通常是大规模数据并行处理的理想select,这是GPU所擅长的。

如果是的话:有谁知道任何教程或得到任何示例代码(编程语言无所谓)?

  • 当你想要去GPGPU的方式,你有两个select: CUDAOpenCL

    CUDA成熟了很多工具,但以NVidia GPU为中心。

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

  • 对于教程,您可以通过Rob Farber在CodeProject上获得一系列精彩的文章: http : //www.codeproject.com/Articles/Rob-Farber#Articles

  • 对于您的具体使用情况,有许多用OpenCL绘制直方图的样本(请注意,许多是图像直方图,但原理相同)。

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

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

除了上述海报的build议之外,还可以在适当的情况下使用TPL(任务并行库)在多个内核上并行运行。

上面的例子可以使用Parallel.Foreach和ConcurrentDictionary,但是更复杂的map-reduce安装程序将数组分割成块,每个块生成一个字典,然后将其简化为单个字典,这样可以获得更好的结果。

我不知道你所有的计算是否正确地映射到GPU的function,但是你必须使用map-reducealgorithm来将计算映射到GPU内核,然后把部分结果减less到一个单一的结果,所以你在转移到一个不那么熟悉的平台之前,还可以在CPU上做这件事。

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

我认为把这个单一的stream程应用程序转变成在许多系统上运行的分布式应用程序,并且调整algorithm会大大加快速度,取决于有多less系统可用。

你可以使用经典的“分而治之”的方法。 我将采取的一般方法如下。

使用一个系统将“largeFloatingPointArray”预处理为散列表或数据库。 这将一次完成。 它将使用浮点值作为键,并将数组中出现的次数作为值。 最坏的情况是,每个值只出现一次,但这是不太可能的。 如果largeFloatingPointArray在每次运行应用程序时保持不断变化,那么内存中的哈希表就是有意义的。 如果它是静态的,那么该表可以保存在一个键值数据库(如Berkeley DB)中。 我们称之为“查找”系统。

在另一个系统上,我们称之为“主要”,创build大量工作,并将工作项目分散到N个系统上,并在结果可用时“收集”结果。 例如,工作项目可以像两个数字一样简单,表示系统应该工作的范围。 当一个系统完成工作时,它会发回一系列事件,并准备好处理另一个工作。

性能得到改善,因为我们不会一直迭代largeFloatingPointArray。 如果查找系统成为瓶颈,则可以根据需要在多个系统上复制。

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

我正在开发针对多核系统(通常称为微服务器)的针对C语言并行编程的编译器,这些系统将在系统中使用多个“片上系统”模块构build。 ARM模块供应商包括Calxeda,AMD,AMCC等。英特尔也可能会有类似的产品。

我有一个版本的编译器工作,可以用于这样的应用程序。 基于C函数原型的编译器生成在整个系统中实现进程间通信代码(IPC)的Cnetworking代码。 其中一种可用的IPC机制是socket / tcp / ip。

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

2012年11月16日添加。

我想了一点关于这个algorithm,我认为这应该在一个单一的通行证。 它是用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. * * Eg 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]++; } } }