更新 GPU 版本
__global__ void hash (float *largeFloatingPointArray,int largeFloatingPointArraySize, int *dictionary, int size, int num_blocks)
{
int x = (threadIdx.x + blockIdx.x * blockDim.x);
float y;
int noOfOccurrences = 0;
int a;
while( x < size )
{
dictionary[x] = 0;
noOfOccurrences = 0;
for(int j = 0 ;j < largeFloatingPointArraySize; j ++)
{
y = largeFloatingPointArray[j];
y *= 10000;
a = y + 0.5;
if (a == x) noOfOccurrences++;
}
dictionary[x] += noOfOccurrences;
x += blockDim.x * gridDim.x;
}
}
这个我只是测试了一下小的输入,因为我是在我的笔记本电脑上进行测试。尽管如此,它仍然可以工作,但需要进行更多测试。
更新 顺序版本
我刚刚做了这个天真的版本,它可以在不到20秒的时间内执行包含30000000个元素的数组的算法(包括生成数据的函数所花费的时间)。
这个天真的版本首先对浮点数数组进行排序。之后,将遍历排序后的数组,并检查给定的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));
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;
}
如果您的计算机安装了thrust库,则应使用以下内容:
#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(
{
}
}
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中,您基本上需要将每个
for迭代分配给每个线程。
1) threadIdx.x + blockIdx.x*blockDim.x;
每个块都有一个ID,从0到N-1(N是块的最大数量),每个块都有一个“X”线程数,每个线程都有一个ID,从0到X-1。给出了for循环迭代,每个线程将根据其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)
int x = blockIdx.x;
float y;
while( x < 1000000)
{
int noOfOccurrences = 0;
float z = converting_int_to_float(x);
for(j = threadIdx.x; j < largeFloatingPointArraySize; j += blockDim.x)
{
y = largeFloatingPointArray[j];
if (z == y)
{
noOfOccurrences++;
}
}
if(threadIdx.x == 0)
atomicAdd(&dictionary[x], noOfOccurrences);
__syncthreads();
}
你需要使用atomicAdd
,因为来自不同块的不同线程可能会同时读/写noOfOccurrences
,所以你必须确保互斥性。
这只是一种方法;你甚至可以将外部循环的迭代分配给线程而不是块。
教程
Rob Farmer撰写的《CUDA:超级计算》系列文章在其十四个部分中涵盖了几乎所有内容,非常适合初学者入门。
还有其他资源:
看最后一个项目,你会发现许多学习CUDA的链接。
OpenCL:OpenCL教程 | MacResearch