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
}
}
これは、ラップトップでテストしているため、小さな入力用にテストしたばかりです。それにもかかわらず、それは機能していますが、さらにテストが必要です。
UPDATE順次バージョン
30,000,000 要素を持つ配列のアルゴリズムを 20 秒未満で実行するこの単純なバージョンを実行しました (データを生成する関数にかかる時間を含む)。
この単純なバージョンでは、最初に float の配列を並べ替えます。その後、ソートされた配列を調べて、指定されたものが配列に出現する回数をチェックし、valueこの値を出現回数とともに辞書に入れます。
私が使っsortedた の代わりに mapを使うことができます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);
確かに速くなります。
元の投稿
私は、1,000 万から 3,000 万の浮動小数点値を含む大きな配列を持つ統計アプリケーションに取り組んでいます。
GPU を利用してそのような計算を高速化することは可能ですか (それは理にかなっていますか)?
はい、そうです。1 か月前、GPU 上で完全な分子動力学シミュレーションを実行しました。粒子のペア間の力を計算したカーネルの 1 つは、パラメーター6配列としてそれぞれ500,000double を受け取り、合計で3100 万 doubleになり(22 MB)ました。
したがって、グローバル メモリの30約 100 万個の浮動小数点数を配置する予定がある場合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 のコードの簡単な例は、2 つの 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 では、基本的にそれぞれを繰り返し処理し、各スレッドに割り当てます。
1) threadIdx.x + blockIdx.x*blockDim.x;
各ブロックにはIDfrom0からN-1(N はブロックの最大数) があり、各ブロックにはfrom からまで'X'のスレッドがいくつかあります。ID0X-1
- 各スレッドがそのスレッドとスレッドが含まれるブロックに基づいて計算する for ループの反復を提供します。blockDim.x は、ブロックが持つスレッドの数です。
IDID
したがって、それぞれに10スレッド とがある 2 つのブロックがある場合、次のようになります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同時に書き込み/読み取りを行う可能性があるため、使用する必要があるため、相互排除を確保する必要があります。
これは 1 つのアプローチにすぎません。外側のループの反復をブロックではなくスレッドに割り当てることもできます。
チュートリアル
Dr Dobbs Journal シリーズのCUDA: Rob Farmer による大衆向けのスーパーコンピューティングは優れており、14 回の記事でほぼすべてをカバーしています。それはまたかなり穏やかに始まるので、かなり初心者に優しい.
その他:
最後の項目を見ると、CUDA を学ぶためのリンクがたくさんあります。
OpenCL: OpenCL チュートリアル | マックリサーチ