GPU 如何帮助改进迭代问题

How does GPU help in improving iterative questions?

本文关键字:迭代 问题 帮助 何帮助 GPU      更新时间:2023-10-16

我正在使用C++求解k耦合迭代方程。例如,对于 3 联轴器情况,

 f(n+1) = g(n) + 2*h(n) + c;
 g(n+1) = 0.5*f(n+1) - h(n);
 h(n+1) = ( f(n+1)+g(n+1) )/2; 

其中 C 是常量。在 C/C++ 中,实现非常简单

#include <vector>
#include <iostream>
using namespace std;
void main(void)
{
  double c= 0.24;
  long k=0;
  vector<double> f(900000000), g(900000000), h(900000000);
  while (k<10000000)
  {
    f[0] = g[0] = h[0] = rand(); // the initial values of f, g, h are randomly picked
    for (long n=1; n<900000000; n++)
    {
      f[n+1] = g[n] + 2*h[n] + c;
      g[n+1] = 0.5*f[n+1] - h[n];
      h[n+1] = ( f[n+1]+g[n+1] )/2; 
    }
    //if the final value of f, g, h satisfying some condition then record it and go for next iteration 
    if (is_good(f[899999999], g[899999999], h[899999999]))
    {
      // record f[899999999], g[899999999], h[899999999]
      k++;
    }
  }
}

这段代码非常慢,因为它进展缓慢并且依赖于随机首字母。我以前没有编写过 GPU 代码,但我读了一些介绍,它说 GPU 在某些情况下非常快。我读了一些例子,我觉得GPU只能用于"可分割"的情况(我的意思是任务可以划分为子任务,因此可以并行实现(。我想知道这对我的案子有多大帮助。任何想法或建议将受到高度欢迎。

您的程序可以在while (k<10000000)循环中轻松并行化。 事实上,由于程序终止条件是未知的迭代次数(达到 10M 个好集合(,你基本上可以删除你在内核中显示的整个代码并按原样运行它,只需进行一些小的修改。

#include <curand.h>
#include <curand_kernel.h>
__constant__ double c = 0.24;
__device__ volatile unsigned int k = 0;
#define SCALE 32767.0
#define NUM_GOOD 10000000
__device__ int is_good(double f, double g, double h){
  if (....){
    ...
    return 1;
  }
  return 0;
}
__global__ void initCurand(curandState *state, unsigned long seed){
  int idx = threadIdx.x + blockIdx.x*blockDim.x;
  curand_init(seed, idx, 0, &state[idx]);
}
__global__ void mykernel(curandState *devStates, double *good_f, double *good_g, double *good_h){
  int idx = threadIdx.x + blockDim.x*blockIdx.x;
  double f0, g0, h0, f1, g1, h1;
  curandState localState = devStates[idx];
  while (k<NUM_GOOD){
    // assuming you wanted independent starting values for f, g, h
    f0 = (double)(curand_uniform(&localState)*SCALE);
    g0 = (double)(curand_uniform(&localState)*SCALE);
    h0 = (double)(curand_uniform(&localState)*SCALE);
    for (int i = 0; i< 450000000; i++){
      f1 = g0 + 2*h0 + c;
      g1 = 0.5*f1 - h0;
      h1 = (f1+g1 )/2;
      f0 = g1 + 2*h1 + c;
      g0 = 0.5*f0 - h1;
      h0 = (f0+g0 )/2;}
    if (is_good(f1, g1, h1))
    {
      unsigned int next =  atomicAdd(&k, 1);
      if (next<NUM_GOOD){
        good_f[next] = f1;
        good_g[next] = g1;
        good_h[next] = h1;}
    }
  }
}

上面的代码只是一个大纲,可能存在一些错误,显然不是这里定义了所有内容。

您可以使用您启动的实际线程数来查看运行最快的线程。 所有启动的线程都将填充"好"堆栈,直到它被填充。 然后每个线程将检测到堆栈已满并退出。

编辑:回答以下一些问题:

似乎"int idx = threadIdx.x + blockDim.x*blockIdx.x;"是 GPU 的东西,我认为它与 GPU 中的线程有关,那么它对 GPU 编程至关重要吗?

是的,这些变量(如 threadIdx.x 是 CUDA 中的"内置"变量,它允许每个线程执行不同的事情(在这种情况下,从不同的随机值开始(。

其次,您提供的所有代码看起来都像常规C++代码。但是你放了"GPU 关键部分",那么我需要在该部分使用任何特殊语法还是就像常规的 c++ 代码一样?

是的,大部分 CUDA 内核代码可以是普通的C++代码,通常类似于您在 CPU 上编写的相同操作的代码。 在这种情况下,我提到了一个关键部分并链接了一个示例,但是在考虑之后,关键部分(在这种情况下用于限制对数据区域的访问,以便 GPU 线程在更新"良好"值时不会相互踩踏(在这里是矫枉过正的。 只需要使用原子操作在堆栈中为每个想要填充良好值的线程保留一个"点"。 我已经相应地修改了代码。

根据

 while (k<10000000)

您正在尝试找到 10M 的好{f, h, g}.

在你的单线程CPU代码中,你一个接一个地找到它们,而在GPU中,很容易启动数千个线程并行找到满意的结果,直到总数达到10M。

对于耦合迭代部分,您仍然需要以传统方式计算它们。但是您仍然能够通过简化方程来提高这部分的性能,

例如
f(n+1) = 1   *g(n) + 2*h(n) +      c;
g(n+1) = 0.5 *g(n)          +  0.5*c;
h(n+1) = 0.75*g(n) + 1*h(n) + 0.75*c;

向量[f,g,h,c]'的变换矩阵A为(在 matlab 代码中(

A = [ 0 1 2 1 ; 0 .5 0 .5; 0 .75 1 .75 ; 0 0 0 0];

然后我们有[f,g,h,c]'{n}=A^n * [f,g,h,c]'{0}.您会发现A^n在几次迭代中收敛到[0 3 2 3; 0 0 0 0; 0 1.5 1 1.5; 0 0 0 0]