多GPU的实现问题

  • 6 replies
  • 202 views
多GPU的实现问题
« 于: 一月 14, 2020, 04:36:44 pm »
新手,想在一个原有单GPU的系统上实现多GPU版本
我在网上看到的多GPU实现都是只有一个内核函数,举例代码如下:
#include <iostream>
#include <stdlib.h>
#include <math.h>
#include <cuda_runtime.h>
// This application demonstrates how to use CUDA API to use mutiple GPUs(4 Nvidia Geforce GTX 1080 ti)
// Function to add the elements of two arrays
 
//Mutiple-GPU Plan Structure
typedef struct
{
    //Host-side input data
    float *h_x, *h_y;
    
    //Result copied back from GPU
     float *h_yp;
    //Device buffers
    float *d_x, *d_y;
 
    //Stream for asynchronous command execution
    cudaStream_t stream;
 
} TGPUplan;
 
// CUDA Kernel function to add the elements of two arrays on the GPU
__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x + blockIdx.x * blockDim.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
      y = x + y;
}
 
int main(void)
{
  int N = 1<<20; // 1M elements
 
  //Get the numble of CUDA-capble GPU
  int N_GPU;
  cudaGetDeviceCount(&N_GPU);
 
  //Arrange the task of each GPU
  int Np = (N + N_GPU - 1) / N_GPU;
 
  //Create GPU plans
  TGPUplan plan[N_GPU];
 
  //Initializing
  for(int i = 0; i < N_GPU; i++)
  {
    cudaSetDevice(i);
    cudaStreamCreate(&plan.stream);
 
    cudaMalloc((void **)&plan.d_x, Np * sizeof(float));
    cudaMalloc((void **)&plan.d_y, Np * sizeof(float));
    plan.h_x = (float *)malloc(Np * sizeof(float));
    plan.h_y = (float *)malloc(Np * sizeof(float));
    plan.h_yp = (float *)malloc(Np * sizeof(float));
 
     for(int j = 0; j < Np; j++)
    {
      plan.h_x[j] = 1.0f;
      plan.h_y[j] = 2.0f;
    }
  }
 
  int blockSize = 256;
  int numBlock = (Np + blockSize - 1) / blockSize;
 
  for(int i = 0; i < N_GPU; i++)
  {
    //Set device
    cudaSetDevice(i);
 
    //Copy input data from CPU
    cudaMemcpyAsync(plan.d_x, plan.h_x, Np * sizeof(float), cudaMemcpyHostToDevice, plan.stream);
    cudaMemcpyAsync(plan.d_y, plan.h_y, Np * sizeof(float), cudaMemcpyHostToDevice, plan.stream);
   
    //Run the kernel function on GPU
    add<<<numBlock, blockSize, 0, plan.stream>>>(Np, plan.d_x, plan.d_y);
   
    //Read back GPU results
    cudaMemcpyAsync(plan.h_yp, plan.d_y, Np * sizeof(float), cudaMemcpyDeviceToHost, plan.stream);
  }
 
  //Process GPU results
  float y[N];
  for(int i = 0; i < N_GPU; i++)
  {
    //Set device
    cudaSetDevice(i);
 
    //Wait for all operations to finish
    cudaStreamSynchronize(plan.stream);
 
    //Get the final results
     for(int j = 0; j < Np; j++)
        if(Np * i + j < N)
            y[Np * i + j]=plan.h_yp[j];
    
    //shut down this GPU
    cudaFree(plan.d_x);
    cudaFree(plan.d_y);
    free(plan.h_x);
    free(plan.h_y);
     cudaStreamDestroy(plan.stream); //Destroy the stream
  }
 
  // Check for errors (all values should be 3.0f)
  float maxError = 0.0f;
  for (int i = 0; i < N; i++)
    maxError = fmax(maxError, fabs(y-3.0f));
  std::cout << "Max error: " << maxError << std::endl;
 
  return 0;
}
上述代码只是实现了对单个内核函数的多GPU操作,如果想要对多个内核函数进行多GPU操作应该如何进行呢
比如程序有三个内核函数
init<<<    >>>(传参1)
update<<<  >>>(传参2 )
calculate<<<   >>>(传参3)
三个内核函数的传参内容都不相同,这种情况下应该怎样进行编程?

Re: 多GPU的实现问题
« 回复 #1 于: 一月 14, 2020, 06:28:38 pm »
你的init, calculate, update这三个kernel,看上去像是连续的操作,这种情况下,你可以:

(1)将这3个步骤操作的缓冲区,重复申请N份(对应你的卡的数量),例如init可能需要input0, 输出result1; calculate可能需要input1, 和result1, 输出result2; 而update则可能需要input3, result2, 和输出result3。此时将这些缓冲区都重复的申请3份。申请的时候,需要通过cudaSetDevice(你的卡的序号),注意序号从0开始,然后才能开始申请。这样会导致分别申请的缓冲区,对应的到N个卡上。

(2)如果原本1张卡上没有使用多流,则你只需要建立N个流即可;如果你原本单卡时候已经用了多流了(例如input1的传入,和前一个init的kernel的result1的结果的生成,用单卡多流可以在时间上重叠进行),则你需要原本用的多流数量(例如3个,或者通用的记录成N个),申请3N个流。其中和申请缓冲区一样,先用cudaSetDevice(卡序号), 就能后续的申请的流对应该卡。

(3)反复先用cudaSetDevice()切换到,再反复的切换每个卡上的多个流,以完成你之前的每个卡上的任务处理。

注意这个是1个Host线程,对应多卡的情况。如果你愿意使用多个线程,则情况变得简单:

(A)每个线程中,先cudaSetDevice()切换到该host线程负责的卡上。
(B)后续步骤抄袭原有操作不变。

你可以看到这种方式非常简单,只需要增加Host上开多线程的代码,和设定Device的代码,后续的代码就可以完全不变。我建议你采用这种方式。

Re: 多GPU的实现问题
« 回复 #2 于: 一月 15, 2020, 11:29:15 am »
您好,非常感谢您的回复与帮助。实际上我的三个kernel是分开执行的,大致顺序为
_host_   function()
{
    host code;
    void init();
    host code;
    void update();
    host code;
    void calculate();
    ...
}

其中每个内核函数以嵌套方式在host函数内执行,如下所示
   void update()
{
   locate_leaf_batch();
}
void locate_leaf_batch()
{
  update<<<   >>> ();
}
按照您的方法,我尝试性的用了
1.第一种单Host线程,GPU多流方法
2.多HOST线程方法

对于1方法的单GPU版本如图1,此版本运行无错误

我的修改为多GPU流的部分代码如图2,此版本运行出错

在  语句cudaDeviceSynchronize()上报错,如图3
想请问一下您,我这种做法应该属于单host线程多流吧,出错的原因是什么您可以进行帮助指正么,感谢您的耐心回复。

Re: 多GPU的实现问题
« 回复 #3 于: 一月 15, 2020, 01:41:58 pm »
您好,非常感谢您的回复与帮助。实际上我的三个kernel是分开执行的,大致顺序为
_host_   function()
{
    host code;
    void init();
    host code;
    void update();
    host code;
    void calculate();
    ...
}

其中每个内核函数以嵌套方式在host函数内执行,如下所示
   void update()
{
   locate_leaf_batch();
}
void locate_leaf_batch()
{
  update<<<   >>> ();
}
按照您的方法,我尝试性的用了
1.第一种单Host线程,GPU多流方法
2.多HOST线程方法

对于1方法的单GPU版本如图1,此版本运行无错误

我的修改为多GPU流的部分代码如图2,此版本运行出错

在  语句cudaDeviceSynchronize()上报错,如图3
想请问一下您,我这种做法应该属于单host线程多流吧,出错的原因是什么您可以进行帮助指正么,感谢您的耐心回复。

如果你的kernel可以在原来的缓冲区分配下正常工作的话,而多卡的时候,却报告“访存错误”,则检查你的缓冲区使用。

每个卡上分配的缓冲区,只能在该卡上的分配的流中,所启动的该卡上的kernel所使用。请检查你分配时候设定的设备,和使用时候是否一致,这是最对应你的错误的原因了。

我之前说过,如果你搞不定1个Host线程,对应N个设备上的M个流的话,请直接新建N个Host线程,每个只需要在开头的地方设定Device一次(只调用一次!各个线程就无关了!),后面的代码都不需要动的。但显然你无视了这点,试图上去来个复杂的。

我建议你先学会走路再尝试跑步。

Re: 多GPU的实现问题
« 回复 #4 于: 一月 15, 2020, 01:43:53 pm »
如果你的kernel可以在原来的缓冲区分配下正常工作的话,而多卡的时候,却报告“访存错误”,则检查你的缓冲区使用。

每个卡上分配的缓冲区,只能在该卡上的分配的流中,所启动的该卡上的kernel所使用。请检查你分配时候设定的设备,和使用时候是否一致,这是最对应你的错误的原因了。

我之前说过,如果你搞不定1个Host线程,对应N个设备上的M个流的话,请直接新建N个Host线程,每个只需要在开头的地方设定Device一次(只调用一次!各个线程就无关了!),后面的代码都不需要动的。但显然你无视了这点,试图上去来个复杂的。

我建议你先学会走路再尝试跑步。

此外,你文字说的是,你使用了”多CPU线程的版本“,然而很遗憾,从代码的截图看,你并没有这样。

(你的代码在反复的切换设备/流,这就属于刚才说的跑的方式了。请先使用简单的走路的方式。 CPU上的多线程可以有效的简化逻辑的,而不仅仅是为了充分利用多个核心上的硬件线程,请一定要相信这点)。

Re: 多GPU的实现问题
« 回复 #5 于: 一月 16, 2020, 11:58:08 am »
您好,非常感谢您的耐心回复
我在之前发此贴过程中正在做多CPU线程的方法,您的建议对我非常有帮助,现在我的多CPU线程方法出现了 cudaErrorIllegalAddress问题,这是我的代码

如代码所示,我现在想要将输入的数据gpma.keys分为两个部分,通过data.p_tree_size=gpma.keys.size()   
分别给两个线程并行处理。 在运行过程中出现了代码错误如下图,您能根据代码看出我的越界错误出在哪里么,非常感谢您的帮助

Re: 多GPU的实现问题
« 回复 #6 于: 一月 16, 2020, 08:56:32 pm »
抱歉,不懂Thrust和C++的,暂时无法为这两个方面提供技术支持。