两个CPU线程通过不同的对象分别调用同一个GPU进行处理,是否可行?

  • 7 replies
  • 205 views
求助问题:在CPU端通过OpenMP开了两个线程,每个线程在同一GPU设备上分别创建一个stream, 调用同一个功能函数(内含若干个kernel函数),两个线程均会调度失败。每个线程分别单独执行,均能成功。每个线程的资源以句柄的方式开辟,互不干扰。暂时不知道哪里出问题了?

求助问题:在CPU端通过OpenMP开了两个线程,每个线程在同一GPU设备上分别创建一个stream, 调用同一个功能函数(内含若干个kernel函数),两个线程均会调度失败。每个线程分别单独执行,均能成功。每个线程的资源以句柄的方式开辟,互不干扰。暂时不知道哪里出问题了?

OpenMP我不清楚,但通过常规的手段创建两个CPU线程的话,这样使用是无问题的。

不过这种无问题是针对CUDA Runtime API本身说的,如果你还有一层“功能函数”进行包装,这里是否多线程安全,则要看你自己的代码了。

CUDA本身这样做是毫无问题的。

谢谢您的解答,问题确实如您的回答,是可以通过两个CPU线程去调度同一个GPU进行处理的。
我通过排查错误发现是kernel函数没有启动成功。

今天我在K80上跑这个简单的求向量的和的demo时,在启动配置不变时,增加数据长度至100M时,kernel会失败,因此,失败的原因不是双线程调度同一个GPU的问题。
因此,想再请教个问题,如这样<<< >>>>启动配置中,配置的总的线程数是不是不能任意大呢?是否有一个限制呢?

下面为测试小demo:
1->test.h
#pragma once

#pragma once

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "omp.h"

#include <stdio.h>
#include <stdlib.h>
#include <windows.h>

struct xyhandle{
   cudaStream_t stream;
   int *h_array_A;
   int *d_array_A;
   int *h_array_B;
   int *d_array_B;
   int *h_array_C;
   int *d_array_C;
};

int init_function(xyhandle *pHandle, int dataLen);
void pro_function(xyhandle *pHandle, int dataLen, int initValue);

2->kernel.cu
#include "test.h"

__global__ void addKernel(int *c, const int *a, const int *b, int dataLen)
{
   int tid = threadIdx.x + blockIdx.x * blockDim.x;//具体位置的索引

   while (tid < dataLen)
   {
      c[tid] = a[tid] + b[tid];
      tid += blockDim.x * gridDim.x;
   }
}



int init_function(xyhandle *pHandle, int dataLen)
{
   cudaSetDevice(0);
   cudaMallocHost(&(pHandle->h_array_A), dataLen*sizeof(int));
   cudaMallocHost(&(pHandle->h_array_B), dataLen*sizeof(int));
   cudaMallocHost(&(pHandle->h_array_C), dataLen*sizeof(int));
   if ((pHandle->h_array_A == NULL) || (pHandle->h_array_B == NULL)||(pHandle->h_array_C==NULL))
   {
      return -1;
   }
   cudaMalloc(&(pHandle->d_array_A), dataLen*sizeof(int));
   cudaMalloc(&(pHandle->d_array_B), dataLen*sizeof(int));
   cudaMalloc(&(pHandle->d_array_C), dataLen*sizeof(int));
   if ((pHandle->d_array_A == NULL) || (pHandle->d_array_B == NULL) || (pHandle->d_array_C==NULL))
   {
      return -1;
   }

   cudaStreamCreate(&(pHandle->stream));
   if (pHandle->stream == NULL)
   {
      return -1;
   }

   return 0;
}

void pro_function(xyhandle *pHandle, int dataLen, int initValue)
{
   cudaError_t cudaStatus;
   cudaStream_t stream = pHandle->stream;
   int *h_array_A = pHandle->h_array_A;
   int *h_array_B = pHandle->h_array_B;
   int *h_array_C = pHandle->h_array_C;
   int *d_array_A = pHandle->d_array_A;
   int *d_array_B = pHandle->d_array_B;
   int *d_array_C = pHandle->d_array_C;

   // 初始化
   for (int i = 0; i < dataLen; ++i)
   {
      h_array_A = i * 2 + initValue;
      h_array_B = i * 1 + initValue;
   }


   cudaMemcpyAsync(d_array_A, h_array_A, dataLen*sizeof(int), cudaMemcpyHostToDevice, stream);
   cudaMemcpyAsync(d_array_B, h_array_B, dataLen*sizeof(int), cudaMemcpyHostToDevice, stream);
   
   int threadNum = 128;
   int blockNum  = (dataLen + threadNum - 1) / threadNum;
   addKernel << <blockNum, threadNum, 0, stream >> >(d_array_C, d_array_A, d_array_B, dataLen);
   
   cudaMemcpyAsync(h_array_C, d_array_C, dataLen*sizeof(int), cudaMemcpyDeviceToHost, stream);
   
   cudaStatus = cudaStreamSynchronize(stream);
   if (cudaStatus != cudaSuccess)
   {
      printf("同步失败!\n");
   }
   printf("run here!\n");
   // 验证
   for (int i = 0; i < dataLen; ++i)
   {
      if (h_array_A + h_array_B != h_array_C)
      {
         printf("执行失败!\n");
         return;
      }
   }

   printf("执行成功!\n");
}


3->test.cpp
#include "test.h"


int main()
{
   int dataLen = 100000000;
   int initValue = 22;
   xyhandle Handle[2];
   if (init_function(&(Handle[0]), dataLen) != 0)
   {
      printf("初始化失败!\n");
      return -1;
   }
   if (init_function(&(Handle[1]), dataLen) != 0)
   {
      printf("初始化失败!\n");
      return -1;
   }


#pragma omp parallel for
   for (int i = 0; i < 2; ++i)
   {
      printf("第%d次执行!\n", i);
      pro_function(&(Handle), dataLen, initValue);
   }


   system("pause");

   return 0;
}

« 最后编辑时间: 十一月 08, 2019, 08:45:50 pm 作者 hoptony »

谢谢您的解答,问题确实如您的回答,是可以通过两个CPU线程去调度同一个GPU进行处理的。
我通过排查错误发现是kernel函数没有启动成功。

今天我在K80上跑这个简单的求向量的和的demo时,在启动配置不变时,增加数据长度至100M时,kernel会失败,因此,失败的原因不是双线程调度同一个GPU的问题。
因此,想再请教个问题,如这样<<< >>>>启动配置中,配置的总的线程数是不是不能任意大呢?是否有一个限制呢?

下面为测试小demo:
1->test.h
#pragma once

#pragma once

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "omp.h"

#include <stdio.h>
#include <stdlib.h>
#include <windows.h>

struct xyhandle{
   cudaStream_t stream;
   int *h_array_A;
   int *d_array_A;
   int *h_array_B;
   int *d_array_B;
   int *h_array_C;
   int *d_array_C;
};

int init_function(xyhandle *pHandle, int dataLen);
void pro_function(xyhandle *pHandle, int dataLen, int initValue);

2->kernel.cu
#include "test.h"

__global__ void addKernel(int *c, const int *a, const int *b, int dataLen)
{
   int tid = threadIdx.x + blockIdx.x * blockDim.x;//具体位置的索引

   while (tid < dataLen)
   {
      c[tid] = a[tid] + b[tid];
      tid += blockDim.x * gridDim.x;
   }
}



int init_function(xyhandle *pHandle, int dataLen)
{
   cudaSetDevice(0);
   cudaMallocHost(&(pHandle->h_array_A), dataLen*sizeof(int));
   cudaMallocHost(&(pHandle->h_array_B), dataLen*sizeof(int));
   cudaMallocHost(&(pHandle->h_array_C), dataLen*sizeof(int));
   if ((pHandle->h_array_A == NULL) || (pHandle->h_array_B == NULL)||(pHandle->h_array_C==NULL))
   {
      return -1;
   }
   cudaMalloc(&(pHandle->d_array_A), dataLen*sizeof(int));
   cudaMalloc(&(pHandle->d_array_B), dataLen*sizeof(int));
   cudaMalloc(&(pHandle->d_array_C), dataLen*sizeof(int));
   if ((pHandle->d_array_A == NULL) || (pHandle->d_array_B == NULL) || (pHandle->d_array_C==NULL))
   {
      return -1;
   }

   cudaStreamCreate(&(pHandle->stream));
   if (pHandle->stream == NULL)
   {
      return -1;
   }

   return 0;
}

void pro_function(xyhandle *pHandle, int dataLen, int initValue)
{
   cudaError_t cudaStatus;
   cudaStream_t stream = pHandle->stream;
   int *h_array_A = pHandle->h_array_A;
   int *h_array_B = pHandle->h_array_B;
   int *h_array_C = pHandle->h_array_C;
   int *d_array_A = pHandle->d_array_A;
   int *d_array_B = pHandle->d_array_B;
   int *d_array_C = pHandle->d_array_C;

   // 初始化
   for (int i = 0; i < dataLen; ++i)
   {
      h_array_A = i * 2 + initValue;
      h_array_B = i * 1 + initValue;
   }


   cudaMemcpyAsync(d_array_A, h_array_A, dataLen*sizeof(int), cudaMemcpyHostToDevice, stream);
   cudaMemcpyAsync(d_array_B, h_array_B, dataLen*sizeof(int), cudaMemcpyHostToDevice, stream);
   
   int threadNum = 128;
   int blockNum  = (dataLen + threadNum - 1) / threadNum;
   addKernel << <blockNum, threadNum, 0, stream >> >(d_array_C, d_array_A, d_array_B, dataLen);
   
   cudaMemcpyAsync(h_array_C, d_array_C, dataLen*sizeof(int), cudaMemcpyDeviceToHost, stream);
   
   cudaStatus = cudaStreamSynchronize(stream);
   if (cudaStatus != cudaSuccess)
   {
      printf("同步失败!\n");
   }
   printf("run here!\n");
   // 验证
   for (int i = 0; i < dataLen; ++i)
   {
      if (h_array_A + h_array_B != h_array_C)
      {
         printf("执行失败!\n");
         return;
      }
   }

   printf("执行成功!\n");
}


3->test.cpp
#include "test.h"


int main()
{
   int dataLen = 100000000;
   int initValue = 22;
   xyhandle Handle[2];
   if (init_function(&(Handle[0]), dataLen) != 0)
   {
      printf("初始化失败!\n");
      return -1;
   }
   if (init_function(&(Handle[1]), dataLen) != 0)
   {
      printf("初始化失败!\n");
      return -1;
   }


#pragma omp parallel for
   for (int i = 0; i < 2; ++i)
   {
      printf("第%d次执行!\n", i);
      pro_function(&(Handle), dataLen, initValue);
   }


   system("pause");

   return 0;
}

hoptony,你这代码完全没有进行任何错误检查,包括每次内存分配和显存分配(这往往是大头)。注意你检查是否为NULL是不安全的(正常的做法是检查返回值是否为cudaSuccess, 如果非要这样检查,请在分配前设定指针为NULL)。请先所有的CUDA返回值。

此外,100M(元素单位)不大,常规的blocks至少可以到2^31-1(x方向),你如果是100M个总线程,然后每个block才128个线程,这个才800K,不是这个问题的。

暂时无更多结论,请先按照规范写法写(检查返回值),往往此时错误可以自然暴露。


hoptony,你这代码完全没有进行任何错误检查,包括每次内存分配和显存分配(这往往是大头)。注意你检查是否为NULL是不安全的(正常的做法是检查返回值是否为cudaSuccess, 如果非要这样检查,请在分配前设定指针为NULL)。请先所有的CUDA返回值。

此外,100M(元素单位)不大,常规的blocks至少可以到2^31-1(x方向),你如果是100M个总线程,然后每个block才128个线程,这个才800K,不是这个问题的。

暂时无更多结论,请先按照规范写法写(检查返回值),往往此时错误可以自然暴露。

此外,注意你的排版,至少一处或者多处下标消失了(请使用代码模式)。
以及,任何涉及到大小的地方均建议使用size_t,  哪怕你这里用普通的int足够了。

谢谢您的耐心指导和建议,实际工作中的相关模块,已正常走通。                        确实测试demo写得不规范,也不严谨,以后一定改正,做技术确实是要很严谨。
之前提的问题,对于测试demo,在k80和cuda 7.5的环境下,经排查,启动配置线程块数过大时,确实kernel没启动起来,减少线程块数量,在相同数据量的情况下,kernel能正常起来,此问题留待后面排查。
« 最后编辑时间: 十一月 11, 2019, 10:29:27 pm 作者 hoptony »

谢谢您的耐心指导和建议,实际工作中的相关模块,已正常走通。                        确实测试demo写得不规范,也不严谨,以后一定改正,做技术确实是要很严谨。
之前提的问题,对于测试demo,在k80和cuda 7.5的环境下,经排查,启动配置线程块数过大时,确实kernel没启动起来,减少线程块数量,在相同数据量的情况下,kernel能正常起来,此问题留待后面排查。

我司从计算能力1.x时代开始跟踪NVIDIA GPU,没有发现有显卡不能达到手册中标称的grid规模的(即,你这里的blocks的数量,特别是X方向),这10年来未有一例用户报告过grid规模不能达标的问题。

根据一般规律,你的分析结论很可能是错误的。建议重新排查。

我司从计算能力1.x时代开始跟踪NVIDIA GPU,没有发现有显卡不能达到手册中标称的grid规模的(即,你这里的blocks的数量,特别是X方向),这10年来未有一例用户报告过grid规模不能达标的问题。

根据一般规律,你的分析结论很可能是错误的。建议重新排查。

如果之前我们的回复指出的,不到1M的blocks总数,实际上一点都不大,真心不存在这种情况的blocks数量下,kernel不能启动的问题。你应当检查的是之前的所有CUDA API调用的错误代码(然而你并没有这样做)。靠猜测是得不到结论的。

本主题到此结束。