CUDA学习笔记 —— (五)CPU于GPU并行计算实例对比

CUDA学习笔记 —— (五)CPU于GPU并行计算实例对比,第1张

文章目录
      • 概述
      • CPU运行
      • GPU运行单线程
      • GPU 1block 多线程运行
      • 多block多线程
      • 结果比对

概述

我们通过一个例子,相同的计算,我们分别在CPU,GPU单线程,GPU单block多线程,GPU多block多线程来运行对比。


看GPU是如何大幅提升运算能力的。



参考文章https://developer.nvidia.com/blog/even-easier-introduction-cuda/

CPU运行

一个最简单的例子:将两个数组中的元素相加放到第二个数组中。



我们使用CPU来运行

#include 
#include 

// function to add the elements of two arrays
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
      y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20; // 1M elements

  float *x = new float[N];
  float *y = new float[N];

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the CPU
  add(N, x, y);

  // 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[i]-3.0f));
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory
  delete [] x;
  delete [] y;

  return 0;
}
GPU运行单线程

下面作为对比,我们看如何将计算放入到GPU中来完成
下面注释中有具体说明每个函数使用情况。


#include 
#include 
// Kernel function to add the elements of two arrays
__global__
void add(int n, float *x, float *y)
{
  for (int i = 0; i < n; i++)
    y[i] = x[i] + y[i];
}

int main(void)
{
  int N = 1<<20;
  float *x, *y;

  // Allocate Unified Memory – accessible from CPU or GPU
  //这个Unified Memory内存空间使用cudaMallocManaged创建,该内存可以在CPU和GPU之间共享
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));

  // initialize x and y arrays on the host
  for (int i = 0; i < N; i++) {
    x[i] = 1.0f;
    y[i] = 2.0f;
  }

  // Run kernel on 1M elements on the GPU
  //使用1block,1个thread来执行kernel
  add<<<1, 1>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  //等GPU执行完毕,写会内存,因为上面函数是异步过程,所以如果不等GPU写完则会出现数据竞争情况。


cudaDeviceSynchronize(); // 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[i]-3.0f)); std::cout << "Max error: " << maxError << std::endl; // Free memory cudaFree(x); cudaFree(y); return 0; }

GPU 1block 多线程运行

我们修改运行的线程数量

add<<<1, 256>>>(N, x, y);

然后修改kernel函数
因为线程数量为256个,总共有n个数需要计算
这里将N个数平均分到256个线程去运行

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x; //0
  int stride = blockDim.x; //线程数
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}
多block多线程

因为GPU有多个并行的处理器,叫做SM。


而每个SM可以并发运行blocks。



所以我们开启多个block来最大化GPU运行能力。


int blockSize = 256; //开启256个线程
int numBlocks = (N + blockSize - 1) / blockSize; //block数为N/256,但是要注意四舍五入
add<<<numBlocks, blockSize>>>(N, x, y);

我们可以调用变量来获取线程id,一维对应关系如下

还需要改造一下kernel函数,让每个线程对应于一个数组中的元素

__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x; //线程号
  int stride = blockDim.x * gridDim.x; //总线程数
  //这里用一个for循环来处理,我们叫做grid-stride loop.
  //他是为了即使保证线程数量如果小于元素N数,可以均匀分配
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}
结果比对

下面我们对比一下结果。




可以发现,我们极大的利用了带宽,并且执行时间极大的缩短了。


欢迎分享,转载请注明来源:内存溢出

原文地址: http://outofmemory.cn/langs/563465.html

(0)
打赏 微信扫一扫 微信扫一扫 支付宝扫一扫 支付宝扫一扫
上一篇 2022-04-02
下一篇 2022-04-02

发表评论

登录后才能评论

评论列表(0条)

保存