c++ - 我可以/应该在 GPU 上运行此统计应用程序的代码吗？

Question

我正在开发一个在数组中包含大约 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 来加速此类计算？
• 如果是:有谁知道任何教程或有任何示例代码(编程语言无关紧要)？

Answer 1

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

这个简单的版本首先对你的 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);

肯定会更快。

原帖

I'm working on a statistical application which has a large arraycontaining 10 - 30 millions of floating point values.

Is it possible (and does it make sense) to utilize a GPU to speed upsuch calculations?

是的，是的。一个月前，我在 GPU 上运行了一个完整的分子动力学模拟。其中一个计算粒子对之间的力的内核作为参数接收到 6 数组，每个数组都有 500,000 双倍，总共 3 百万加倍 (22 MB).

所以如果你打算放30百万个 float ，也就是大约114MB的全局内存，那是没有问题的。

在您的情况下，计算次数会成为问题吗？根据我在分子动力学 (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 中，您基本上将每次 for 迭代并分配给每个线程，

1) threadIdx.x + blockIdx.x*blockDim.x;

每个 block 都有一个从0到N-1的ID(N是 block 的最大数量)每个 block 都有一个'X' ID 从 0 到 X-1 的线程数。

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

因此，如果您有 2 个 block ，每个 block 具有 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 因为来自不同 block 的不同线程可能同时写入/读取 noOfOccurrences，因此您必须确保 mutual exclusion .

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

教程

Dobbs 博士期刊系列 CUDA: Supercomputing for the masses Rob Farmer 的作品非常出色，在其十四期中几乎涵盖了所有内容。它的启动也相当温和，因此对初学者非常友好。

还有其他人:

• Volume I: Introduction to CUDA Programming
• Getting started with CUDA
• CUDA Resources List

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

OpenCL:OpenCL Tutorials | MacResearch