第四章 kernel函数基础篇

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();
}