目录
- 前言
- 1. thrust
- 2. error
- 总结
前言
杜老师推出的 tensorRT从零起步高性能部署 课程，之前有看过一遍，但是没有做笔记，很多东西也忘了。这次重新撸一遍，顺便记记笔记。
本次课程学习精简 CUDA 教程-错误处理的理解以及错误的传播特性
课程大纲可看下面的思维导图
1. thrust
thrust 是一个基于 CUDA 的高级并行编程库，提供了一组易于使用的算法和数据结构，用于在 GPU 上进行并行计算。(form chatGPT)
thrust 提供了类似于标准 C++ 模板库（STL）的接口和语法，使开发者能够以一种简单直观的方式利用 GPU 的并行计算能力。它提供了丰富的算法，如排序、归约、扫描、变换等，并支持向量、数组和键值对等多种数据结构。
需要注意的是，thrust 并不是 CUDA 的一部分，而是一个由 NVIDIA 提供的独立库。要使用 thrust，需要将其包含在项目中，并确保正确链接和配置相关的编译器和构建系统。
在后续高性能的开发中，thrust 用的还是较少，但是了解它的存在还是很重要的
thrust 案例的示例代码如下：
#include <stdio.h>
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/sort.h>
#include <iostream>
using namespace std;__host__ __device__
int sort_func(int a, int b){return a > b;
}int main(){int data[] = {5, 3, 1, 5, 2, 0};int ndata = sizeof(data) / sizeof(data[0]);thrust::host_vector<int> array1(data, data + ndata);thrust::sort(array1.begin(), array1.end(), sort_func);thrust::device_vector<int> array2 = thrust::host_vector<int>(data, data + ndata);thrust::sort(array2.begin(), array2.end(), []__device__(int a, int b){return a < b;});printf("array1------------------------\n");for(int i = 0; i < array1.size(); ++i)cout << array1[i] << endl;printf("array2------------------------\n");for(int i = 0; i < array2.size(); ++i)cout << array2[i] << endl;return 0;
}
运行效果如下：
通过 thrust 库，可以轻松地对数据进行排序操作。在这个示例中，thrust::sort() 函数被用来对数据进行排序，用户可以根据自己的需求定义排序函数或使用匿名函数来指定排序规则。这样，开发者可以更加便捷地利用 GPU 的并行计算能力进行高效的数据排序。
关于 thrust 的知识点有：(from 杜老师)
- thrust 是 cuda 开发的，基于 cuda 的 stl 库，便于使用
- 因为通常没用到 thrust，所以对这块儿也不做过多解释
- 对于 thrust 中的 lambda 表达式，需要增加 __device__ 标记表明函数可以被核函数调用，此时需要在 Makefile 中增加 --extended-lambda 标记
- 由于使用到了 device vector，因此编译环境需要修改为 nvcc 编译，因此 main.cpp 改成了 main.cu
- 内存的复制和分配已经被 cuda 封装了
2. error
关于 cuda 中的错误处理，你需要了解：
- 错误处理是理解如何控制 cuda 发生的错误和捕获错误的技术
- 在写 cuda 相关代码时，错误检查是错误处理的一种手段
- 在这里着重拿出来讲可恢复和不可恢复的错误，以及其传播的特性
error 案例的示例代码如下：
#include <cuda_runtime.h>
#include <stdio.h>
#include <iostream>
using namespace std;__global__ void func(float* ptr){int pos = blockIdx.x * blockDim.x + threadIdx.x;if(pos == 999){ptr[999] = 5;}
}int main(){float* ptr = nullptr;// 因为核函数是异步的，因此不会立即检查到他是否存在异常func<<<100, 10>>>(ptr);//func<<<100, 1050>>>(ptr);auto code1 = cudaPeekAtLastError();cout << cudaGetErrorString(code1) << endl;// 对当前设备的核函数进行同步，等待执行完毕，可以发现过程是否存在异常auto code2 = cudaDeviceSynchronize();cout << cudaGetErrorString(code2) << endl;// 异常会一直存在，以至于后续的函数都会失败float* new_ptr = nullptr;auto code3 = cudaMalloc(&new_ptr, 100);cout << cudaGetErrorString(code3) << endl;return 0;
}
运行效果如下：
在这个示例中，我们展示了如何处理 CUDA 中的错误。通过调用 cudaPeekAtLastError()，我们可以立即检查最近的 CUDA 错误，而调用 cudaDeviceSynchronize() 则会等待当前设备的核函数执行完毕并检查是否有错误发生。这样，我们可以在程序中适时地处理和报告错误，避免错误在后续的函数调用中传播导致进一步的问题。
关于 error 的知识点有：(from 杜老师)
- 若 cuda 核函数出错，由于他是异步的，立即执行
cudaPeekAtLastError只会拿到对输入参数校验是否正确的状态，而不会拿到核函数是否执行正确的状态- 所以我们最开始是没有捕获到任何异常的
- 也因此需要等待核函数执行完毕后，才真的知道当前核函数是否出错，一般通过设备同步或者流同步进行等待
- 错误可分为可恢复和不可恢复两种：
- 可恢复
- 参数配置错误等，例如 block 越界（一般最大值是1024），shared memory 大小超出范围（一般是48KB）
- 通过
cudaGetLastError可以获取错误代码，同时把当前状态恢复为 success- 该错误在调用核函数后可以立即通过 cudaGetLastError/cudaPeekAtLastError 拿到
- 该错误在下一个函数调用的时候会覆盖
- 不可恢复
- 核函数执行错误，例如访问越界等等异常
- 该错误则会传递到之后的所有 cuda 操作上
- 错误状态通常需要等到核函数执行完毕才能够拿到，也就是有可能在后续的任何流程中突然异常（因为是异步的）
总结
本次课程学习了 thrust 和错误处理，thrust 提供了类似于标准 C++ STL 库的接口和语法，你需要知道有这么一个东西的存在。错误检查在写 cuda 代码时非常重要，对于 cuda 中的错误可分为可恢复和不可恢复两种，在 cuda 程序出现问题时我们需要知道如何处理。
