cuda教程目录

第一章 指针篇
第二章 CUDA原理篇
第三章 CUDA编译器环境配置篇
第四章 kernel函数基础篇
第五章 kernel索引(index)篇
第六章 kenel矩阵计算实战篇
第七章 kenel实战强化篇
第八章 CUDA内存应用与性能优化篇
第九章 CUDA原子(atomic)实战篇
第十章 CUDA流(stream)实战篇
第十一章 CUDA的NMS算子实战篇
第十二章 YOLO的部署实战篇
第十三章 基于CUDA的YOLO部署实战篇

cuda教程背景

随着人工智能的发展与人才的内卷，很多企业已将深度学习算法的C++部署能力作为基本技能之一。面对诸多arm相关且资源有限的设备，往往想更好的提速，满足更高时效性，必将更多类似矩阵相关运算交给CUDA处理。同时，面对市场诸多教程与诸多博客岑子不起的教程或高昂教程费用，使读者(特别是小白)容易迷糊，无法快速入手CUDA编程，实现工程化。
因此，我将结合我的工程实战经验，我将在本专栏实现CUDA系列教程，帮助读者(或小白)实现CUDA工程化，掌握CUDA编程能力。学习我的教程专栏，你将绝对能实现CUDA工程化，完全从环境安装到CUDA核函数编程，从核函数到使用相关内存优化，从内存优化到深度学习算子开发(如：nms)，从算子优化到模型(以yolo系列为基准)部署。最重要的是，我的教程将简单明了直切主题，CUDA理论与实战实例应用，并附相关代码，可直接上手实战。我的想法是掌握必要CUDA相关理论，去除非必须繁杂理论，实现CUDA算法应用开发，待进一步提高，将进一步理解更高深理论。

cuda教程内容

第一章到第三章探索指针在cuda函数中的作用与cuda相关原理及环境配置；

第四章初步探索cuda相关函数编写(global、device、__host__等)，实现简单入门；

第五章探索不同grid与block配置，如何计算kernel函数的index，以便后续通过index实现各种运算；

第六、七章由浅入深探索核函数矩阵计算，深入探索grid、block与thread索引对kernel函数编写作用与影响，并实战多个应用列子(如：kernel函数实现图像颜色空间转换)；

第八章探索cuda内存纹理内存、常量内存、全局内存等分配机制与内存实战应用(附代码)，通过不同内存的使用来优化cuda计算性能；

第九章探索cuda原子(atomic)相关操作，并实战应用(如:获得某些自加索引等)；

第十章探索cuda流stream相关应用，并给出相关实战列子(如:多流操作等)；

第十一到十三章探索基于tensorrt部署yolo算法，我们首先将给出通用tensorrt的yolo算法部署，该部署的前后处理基于C++语言的host端实现，然后给出基于cuda的前后处理的算子核函数编写，最后数据无需在gpu与host间复制操作，实现gpu处理，提升算法性能。

目前，以上为我们的cuda教学全部内容，若后续读者有想了解知识，可留言，我们将根据实际情况，更新相关教学内容。

大神忽略

---

---

前言

本章开始，我们正式进入编程环节。本章介绍cuda编程基础，host或device端如何调用函数，重点说明global、device与host限定词的使用。

---

一、global、device、host的含义

CUDA是通过函数类型的限定词区别函数是否为host或device调用函数，主要以下三个函数类型限定词。

1、global函数

global函数：在device上执行，从host中调用，返回类型必须是void，不支持可变参数，不能成为类成员函数。且__global__修饰的函数用<<<>>>的方式调用，注意用__global__定义的kernel是异步的，这意味着host不会等待kernel执行完就执行下一步。
__global__实际为核函数，后面将有大量使用列子。以下说明核函数形式与参数：

运行时API通过在函数名称和参数列表之间插入<<<Dg, Db, Ns, S>>>的形式来指定。

Dg 的类型为dim3，指定网格的维度和大小，Dg.x * Dg.y 等于所发射的块数量；
Db 的类型为dim3，指定各块的维度和大小，Db.x * Db.y *Db.z 等于各块的线程数量；
Ns 的类型为size_t，指定各块为此调用动态分配的共享存储器（除静态分配的存储器之外），这些动态分配的存储器可供声明为动态数组的其他任何变量使用，Ns 是一个可选参数，默认值为0；
S 的类型为cudaStream t，指定相关流；S 是一个可选参数，默认值为0。

2、device函数

device函数：在device上执行，单仅可以从device中调用，不可以和__global__同时用。

3、host函数

host函数：在host上执行，仅可以从host上调用，一般省略不写，不可以和__global__同时用，但可和__device__同时使用，此时函数会在device和host都编译。

二、host、global、device函数关系

结论：host能调用global函数，global能调用device函数

1、host调用global函数

host调用global函数,类似平常普通函数调用方式，但每个global函数需要<<<Dg, Db, Ns, S>>>参数，代码如下：

test_kernel << <dim3(1), dim3(m*n), 0, nullptr >> > (g_a, g_c);

2、global调用device函数

device是设备上使用的函数，一般只能被global核函数调用，代码如下：

float sigmoid_host(float x) {float y= 1 / (1 + exp(-x));return y;
}
__device__  float sigmoid(float x) {float y= 1 / (1 + exp(-x));return y;
}
__global__ void test_kernel(float* a, float* c) {int idx = threadIdx.x ;c[idx] = sigmoid(a[idx]); //正确方式//c[idx] = sigmoid_host(a[idx]);//绝对错误，无法调用，即：global函数无法调用host函数，只能调用devices函数
}

注意：gloabal 函数绝对无法调用host函数

执行结果如下图：

3、host调用特殊device函数

一般而言，device只能被global函数调用，但有一种特色device函数可被host函数调用，即：函数被host限定词使用，如下sigmod_device_host函数形式，能被host函数调用。具体实现代码如下：

__device__ __host__  float sigmoid_device_host(float x) {float y = 1 / (1 + exp(-x));return y;
}
void host2device(){float y=sigmoid_device_host(1.25);std::cout << y << endl;std::cout << "success：host calling  device+host  " << endl;//以下执行失败   try {float y = sigmoid_host(1.25);throw std::runtime_error("error: fail");   } catch (std::runtime_error err) {std::cout << "fail：host calling device" << endl;}}

执行结果如下：

三、host、global、device函数关系结论

1、函数与设备关系结论

a、host函数无法调用device函数，但可调用__device__ __host__的2个限定函数。

b、device函数在设备gpu上执行，host函数在cpu上执行；

c、global函数通过cpu调用，而global通常为kernel函数，是需要将数据转到gpu上运行。

2、函数间调用形式结论

a、global函数无法调用host函数，可调用device函数；

b、host函数可调用host函数与global函数，可调用组合__device__ __host__函数(实际调用host函数)；

c、device函数可调用device函数；

四、整体代码

函数间调用关系代码如下：
注：附数源码链接[点击这里](https://github.com/tangjunjun966/cuda-tutorial-master)

#include <iostream>
#include <time.h>
#include "opencv2/highgui.hpp"  //实际上在/usr/include下
#include "opencv2/opencv.hpp"
#include "device_launch_parameters.h"
#include <cuda_runtime_api.h>
using namespace cv;
using namespace std;
/*************************************第四节-CUDA函数基础**********************************************/
float sigmoid_host(float x) {float y= 1 / (1 + exp(-x));return y;
}
__device__  float sigmoid(float x) {float y= 1 / (1 + exp(-x));//float y = sigmoid_host(x);return y;
}
__global__ void test_kernel(float* a, float* c) {int idx = threadIdx.x ;c[idx] = sigmoid(a[idx]); //正确方式//c[idx] = sigmoid_host(a[idx]);//绝对错误，无法调用，即：global函数无法调用host函数，只能调用devices函数  
}void Print_dim(float* ptr, int N) {for (int i = 0; i < N; i++){std::cout << "value:\t" << ptr[i] << std::endl;}
}
void init_variables_float(float* a,  int m, int n) {//初始化变量std::cout << "value of a:" << endl;for (int i = 0; i < m; i++) {for (int j = 0; j < n; j++) {a[i * n + j] = rand()/4089 ;std::cout << "\t" << a[i * n + j];}std::cout << "\n";}
}
void global2device() {const int m = 4;const int n = 2;//分配host内存float* a, * c;cudaMallocHost((void**)&a, sizeof(float) * m * n);cudaMallocHost((void**)&c, sizeof(float) * m * n);//变量初始化init_variables_float(a, m, n);// 分配gpu内存并将host值复制到gpu变量中float* g_a;cudaMalloc((void**)&g_a, sizeof(float) * m * n);cudaMemcpy(g_a, a, sizeof(float) * m * n, cudaMemcpyHostToDevice);float* g_c;cudaMalloc((void**)&g_c, sizeof(float) * m * n);test_kernel << <dim3(1), dim3(m * n), 0, nullptr >> > (g_a, g_c);cudaMemcpy(c, g_c, sizeof(float) * m * n, cudaMemcpyDeviceToHost);Print_dim(c, m * n);
}__device__ __host__  float sigmoid_device_host(float x) {float y = 1 / (1 + exp(-x));return y;
}
void host2device(){float y=sigmoid_device_host(1.25);std::cout << y << endl;std::cout << "success：host calling  device+host  " << endl;//以下执行失败   try {float y = sigmoid_host(1.25);throw std::runtime_error("error: fail");   } catch (std::runtime_error err) {std::cout << "fail：host calling device" << endl;}
}void function_criterion_main() {//global2device();//host<--global<--devicehost2device();
}