更新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;
每个块都有一个ID
from 0
to N-1
(N 块的最大数量),每个块都有'X'
多个线程,ID
从0
to X-1
。
- 为您提供每个线程将根据其和线程所在块计算的for循环迭代;blockDim.x 是块拥有的线程数。
ID
ID
因此,如果您有 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 教程 | 麦克研究