cuda编程学习——第二个cuda程序（官方案例分析）！干货向（二）

已于 2023-09-27 20:58:48 修改
阅读量1.7k 收藏 11
点赞数 1
分类专栏： CUDA编程 文章标签： 学习 c++ 开发语言 cuda 新手
于 2023-05-18 19:08:43 首次发布
版权声明：本文为博主原创文章，遵循 CC 4.0 BY-SA 版权协议，转载请附上原文出处链接和本声明。
本文链接：https://blog.csdn.net/qq_40514113/article/details/130733716
版权
本文介绍了如何使用CUDA进行三维重建中的计算加速，从C++模板代码开始，逐步转换为CUDA核函数，并进行内存分配，然后通过并行线程实现GPU上的高效计算，展示了从单线程到多线程并行计算的过程，以提高运算性能。
摘要由CSDN通过智能技术生成

前言：

最近在做三维重建，尤其是Nerf方面多视角合成工作的时候，意识到了cuda的编程计算可以大大提高其中渲染的计算，最明显的例子是Instant-ngp，Plenoxels等文章，因此后面会学Cuda一段时间，同时也就开了这个新坑。
因为笔者也是cuda新手，所以大家有问题的话可以评论区指出，一起学习进步！

我已经更新了我的Github仓库，大家可以前往仓库下载代码
我的CUDA学习仓库
在这里插入图片描述

Nvidia Cuda官方入门资料

运行环境：

Windows10，Visual Studio2019,显卡3050Ti
（大家自行根据自己笔记本情况去配环境，或者没有gpu，租借云服务器去学习也是可以的）

1：C++模板

首先从下面的C++代码开始
其作用为把2个数组元素相加，代码里设为1<<20，也就是在二进制下把00001左移20位，最终值为2^20，为1048576，大约位1million（百万），相加的值最终减去3.0，判断有没有误差。

#include <iostream>
#include <math.h>

// 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));
    //fabs（） 求浮点数的绝对值 fmax()返回2个参数最大的1个
  std::cout << "Max error: " << maxError << std::endl;

  // Free memory

输出如下
在这里插入图片描述

2：编写核函数cuda

（1）__global__可以理解为一个关键词，其告诉编译器该函数可以在GPU上运行，并且可以从CPU代码上调用
（2）这些__global__函数被称为内核，在GPU上运行的代码通常被称为device code，而在CPU上运行的代码被称为host code

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

3：Cuda下的内存分配（Memory Allocation）

（1）为了在GPU上进行计算，我需要分配GPU可访问的内存。CUDA中的统一内存通过提供系统中所有gpu和cpu可访问的单一内存空间使这一点变得容易。
（2）要在统一内存中分配数据，调用cudaMallocManaged()，它返回一个指针，您可以从host(CPU)code或device(GPU)code访问。要释放数据，只需将指针传递给cudaFree()。
（3）我只需要用对cudaMallocManaged()的调用替换上面代码中的new调用，并用对cudaFree（）来执行释放数值操作

  // Allocate Unified Memory -- accessible from CPU or GPU
  float *x, *y;
  cudaMallocManaged(&x, N*sizeof(float));
  cudaMallocManaged(&y, N*sizeof(float));
//理解为，分配空间指向 指针x指向的空间地址，大小为 N* x  x表示float的空间大小，因为有N个浮点数据，大约为1 million

  ...

  // Free memory
  cudaFree(x);
  cudaFree(y);

（4）在GPU上调用核函数需要在参数列表前添加符号<<< >>>
简单理解，启动1个GPU线程来运行add（）函数

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

（5）目前为止，完整代码如下

#include <iostream>
#include <math.h>
// 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
  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
  add<<<1, 1>>>(N, x, y);

  // Wait for GPU to finish before accessing on host
  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;
}

这只是第一步，因为正如所写的那样，这个内核只适用于单个线程，因为运行它的每个线程都会对整个数组执行add操作。此外，由于多个并行线程会读取和写入相同的位置，因此存在竞争条件。

4：并行线程计算

（1）找出内核运行时间的最简单方法是使用nvprof来运行它，nvprof是CUDA工具包附带的命令行GPU分析器。
（2）上述代码我们体验了利用GPU上一个线程来进行计算。那如何并行呢！
下面代码表示CUDA运行时在GPU上启动要使用多少并行线程，这里选取256个线程

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

（3）如果我在运行代码时只做了这样的修改，它将对每个线程执行一次计算，而不是将计算分散到并行线程上。为了正确地做到这一点，我需要修改内核，引入了两个量：
threadIdx.x：包含了the index of the thread within the block，块中的线程索引。此例中，index范围为0~255。
blockDim.x：包含了the size of thread block（number of threads in the thread block）线程块的大小，等于线程块中的线程数量。此例中，该值为256
在这里插入图片描述

__global__
void add(int n, float *x, float *y)
{
  int index = threadIdx.x;
  int stride = blockDim.x;
  for (int i = index; i < n; i += stride)
      y[i] = x[i] + y[i];
}

下面是运行结果（这里其实计算给出结果相较之前会更快）
在这里插入图片描述
官网给出的比较：
这是一个很大的加速(从463ms减少到2.7ms)。
在这里插入图片描述

5：Out of the Blocks

（1）CUDA gpu有许多并行处理器，它们被分组为流多处理器(SMs)。每个SM可以运行多个并发线程块。为了充分利用CUDA GPU，kernel应启动多个thread blocks。

（2）到目前为止，我们可以想到执行配置的第一个参数指定了线程块的数量。这些并行线程块一起构成了所谓的网格grid。因为我有N个元素要处理，每个块有256个线程，所以我只需要计算块的数量来获得至少N个线程。我只是用N除以块大小(注意，如果N不是blockSize的倍数，就要四舍五入)。

int blockSize = 256;//并行线程数量
//使用N 除以块大小，表示 线程块数量
int numBlocks = (N + blockSize - 1) / blockSize;
add<<<numBlocks, blockSize>>>(N, x, y);

（3）我还需要更新内核代码，以考虑到线程块的整个网格。CUDA提供gridDim.x，其中包含网格中的块数量，以及blockIdx.x ,它包含网格中当前线程块的索引。
再引入两个量：
blockIdx.x：包含the index of the block with in the grid。网格中线程块的索引
gridDim.x：包含the size of the grid。网格大小，可以理解为网格中块的数量

图1说明了在CUDA中使用blockDim.x , gridDim.x 和 threadIdx.x对数组(一维)进行索引的方法。其思想是，每个线程通过计算其块开始的偏移量(块索引乘以块大小:blockIdx)来获得其索引。
这个想法是，每个线程通过计算其块开始的偏移量(块索引乘以块大小)来获得其索引:blockIdx.x* blockDim.x)。并在块中添加线程的索引threadIdx.x。
在这里插入图片描述

__global__
void add(int n, float *x, float *y)
{
  int index = blockIdx.x * blockDim.x + threadIdx.x;
  int stride = blockDim.x * gridDim.x;
  for (int i = index; i < n; i += stride)
    y[i] = x[i] + y[i];
}

更新后的kernel还将stride设置为网格中（blockDim.x*gridDim.x）的线程总数。在CUDA内核中，这种类型的循环通常被称为网格跨步循环。

最终代码如下

#include <iostream>
#include <math.h>
#include<stdint.h>
#include<cuda.h>
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
// Kernel function to add the elements of two arrays
__global__
void add(int n, float* x, float* y)
{
    int index = blockIdx.x * blockDim.x + threadIdx.x;
    int stride = blockDim.x * gridDim.x;
    for (int i = index; i < n; i += stride)
        y[i] = x[i] + y[i];
}

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

    // Allocate Unified Memory – accessible from CPU or 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
  //  add <<<1, 256 >>> (N, x, y);


    int blockSize = 256;//并行线程数量
//使用N 除以块大小，表示 线程块数量
    int numBlocks = (N + blockSize - 1) / blockSize;
    add << <numBlocks, blockSize >> > (N, x, y);



    // Wait for GPU to finish before accessing on host
    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;
}

6：结果比较（性能）

在这里插入图片描述

额外知识补充

代码被分成两部分，一部分是在CPU上，也称之为在Host上，另一部分是在GPU上，也称之为在device上。他们两者的关系如下图所示
在这里插入图片描述

网格grids和线程块blocks

网格grids，在上层，至多可以分成三维的blocks，在不同block当中的线程是不能通信的；线程块blocks在相对较低的层级，同样可以将线程分成三维，而在同一个块中的线程是可以通信的。

对于一个核函数，只能有一个grid，但是可以有多个block，之所以将线程划分为grid和block是为了使得结构更清晰，便于线程管理，灵活运用。

在这里插入图片描述

出门吃三碗饭
关注
专栏目录
深度学习运算：CUDA 编程简介
gongdiwudu的专栏
04-26 2579
如今，当我们谈论深度学习时，通常会将其实现与利用 GPU 来提高性能联系起来。GPU（图形处理单元）最初设计用于加速图像、2D 和 3D 图形的渲染。然而，由于它们能够执行许多并行操作，因此它们的实用性超出了深度学习等应用程序。
CUDA入门笔记（二）CUDA编程基础——CPU、GPU的关系
AuntieLee的博客
09-07 1566
参考： 优达学城：https://classroom.udacity.com/courses/cs344/lessons/55120467/concepts/670611900923 CUDA Zone：https://www.nvidia.cn/object/cuda_education_cn_old.html   一、CPU与GPU的关系 1.计算机有两个不同的处理器，CPU（hos...
CUDA教程及其实例程序实现
07-29
CUDA教程及其实例程序实现，非常经典的资料，压缩包中有文档和代码实例，非常好的资料
CUDA案例
Begoniaish的博客
03-31 589
CUDA案例
CUDA 编程
最新发布
撒旦先生的博客
09-25 1016
前几天浅尝了使用 pybind11 将 C++ 代码打包成 Python 包, 假装我已经能通过 C++ 代码加速程序了. 现在, 搞一搞 CUDA 编程, 进而结合 pybind11 把 CUDA 代码打包成 Python 包. 主要目的是了解一下整体流程与框架.
记一次CUDA编程任务
weixin_30689307的博客
02-16 386
　　 　　这个月6号开始，着手解决一个具有实际意义的计算任务。任务数据有9879896条，每条包含30个整数，任务是计算每两条数据之间的斯皮尔相关系数及其P值。原始数据只有500+MB，因此我并不认为这是个多么大的计算任务。随后稍加计算，我还是很惊呆的，要计算(9879896×9879895)÷2≈4.88亿亿组数据，但此时这还只是个数字概念，我也没意识到时间复杂度和空间复杂度的问题。 ...
cuda示例程序
07-15
英伟达官方提供的cuda示例程序，显示cuda设备信息，供cuda入门使用，简单了解cuda原理。
CUDA by example （中文：GPU高性能编程CUDA实战）代码实例
08-28
CUDA by example代码实例
CUDA并行程序设计 GPU编程指南
04-14
第5章介绍与CUDA编程紧密相关的核心概念——网格、线程块与线程，并通过示例说明线程模型与性能的关系。第6章借助实例详细讲解了不同类型内存的工作机制，并指出实践中容易出现的误区。第7章细述多任务的CPU和GPU...
CUDA编程指南5.0版.pdf
09-06
CUDA编程指南 中文版5.0 CUDA编程指南 中文版
CUDA并行程序设计 GPU编程指南-中文英文高清完整版（各500+页）
07-15
CUDA并行程序设计 GPU编程指南-中文英文高清完整版（各500+页）
CUDA 入门小例子，简单实用，上手快，了解CUDA程序的框架
03-09
CUDA 入门小例子，简单实用，上手快，了解CUDA程序的框架，迅速入门
cuda小例子
10-14
通过thrust库对n卡在上层并行。请用nvcc 编译，cu代码同c代码。
CUDA典型例子
11-18
典型cuda编程实例，可以看看，适合cuda编程爱好者
CUDA图像处理示例程序
12-23
一个简单的CUDA程序，使用CUDA进行图像二值化操作，编译环境为VS2005，用到了cuda_vs_wizard进行环境设置，使用OpenCV读入和处理图像并与CUDA的处理结果做比较，这只是一个简单的示例，供入门者学习，程序本身没有做优化。有兴趣的可以做一下优化或在此基础上完成更复杂的图像处理工作
一个成功运行的CUDA小例子
12-24
忙活了很久，终于编译并成功运行了一个自己的CUDA例程，里面有三个小测试程序，包括Hello CUDA！、修改矩阵、矩阵块加
cuda编程实例
learnererer的博客
06-13 2308
no
【CUDA编程】一：从简单的实例入门
qq_43257640的博客
05-31 6095
前提 拥有NVIDIA的显卡- 配置好cuda 有C/C++基础 代码1 //helloworld.cu #include <iostream> #include <stdio.h> /* * __global__ 用于修饰在gpu上运行且由cpu调用的函数 */ __global__ void k_add (int a, int b, int* c ){ *c = a+b; } /* * __host__用于修饰在cpu上运行的函数，一般缺省 */ int main(v
cuda 介绍实例
ly的博客
08-16 1499
编译玩opencv，我们来说说cuda 的基础以及并行计算的基础，写几个测试程序。 我们为什么采用CPU+GPU架构,总的来说可以获取以下的好处： 1）速度快，采用GPU可以获得CPU的数十倍数百倍的性能，可以实时处理高清图像； 2）成本低，在普通PC机器上插上一块1000多元的Geforce GPU卡即可获得数万元、数十万元高端服务器的性能； 3）减少CPU利用率，提高系统的稳定性。 ...
评论 1
添加红包

请填写红包祝福语或标题

个

红包个数最小为10个

元

红包金额最低5元

当前余额3.43元 前往充值 >
需支付：10.00元
成就一亿技术人!
领取后你会自动成为博主和红包主的粉丝 规则
hope_wisdom
发出的红包
实付元
使用余额支付
点击重新获取
扫码支付
钱包余额 0

抵扣说明：

1.余额是钱包充值的虚拟货币，按照1:1的比例进行支付金额的抵扣。
2.余额无法直接购买下载，可以购买VIP、付费专栏及课程。

余额充值