第七章 kenel实战强化篇

前言

学习我的教程专栏，你将绝对能实现CUDA工程化，实现环境安装、index计算、kernel核函数编程、内存优化与steam性能优化、原子操作、nms的cuda算子、yolov5的cuda部署等内容，并开源教程源码。

我相信通过以上学习，已能基本掌握kernel函数的编写，并基本掌握了kernel函数的计算规则。然，有些我们为了巩固以上基础，也为了说明kernel一些额外编码变体。为此，本节我们将通过例子实战kernel的应用。同时，我们也将逐步使用cuda编写深度学习数据预处理的部分kernel函数，但这里只先介绍图像的RGB变成gray图像，后面篇章将会进一步介绍。最终，我们目的是实现yolov5算法加载数据使用cuda编写，yolov5算法后处理也使用cuda编写。

专栏概括

1、cuda教程目录

第一章 指针篇–>点击这里
第二章 CUDA原理篇–>点击这里
第三章 CUDA编译器环境配置篇–>点击这里
第四章 kernel函数基础篇–>点击这里
第五章 kernel索引(index)篇–>点击这里
第六章 kenel矩阵计算实战篇–>点击这里-上篇–点击这里-下篇

第七章 kenel实战强化篇–>点击这里
第八章 CUDA内存应用与性能优化篇–>点击这里-上篇–点击这里-中篇–点击这里-下篇
第九章 CUDA原子(atomic)实战篇–>点击这里
第十章 CUDA流(stream)实战篇–>点击这里
第十一章 CUDA的NMS算子实战篇–>点击这里-上篇–点击这里-下篇
第十二章 YOLO的部署实战篇–>点击这里-上篇–点击这里-下篇

第十三章 基于CUDA的YOLO部署实战篇–>点击这里

2、cuda教程背景

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

3、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教学全部内容，若后续读者有想了解知识，可留言，我们将根据实际情况，更新相关教学内容。

大神忽略

源码链接地址点击这里
yolov5部署代码链接点击这里
yolov5的cuda部署代码链接：文中链接源码

基于我的代码实测，cuda部署yolov5加速10倍，只想说cuda太香了！！！

一、kernel变量打印printf方法

kernel函数的打印，可使用printf函数，如下代码：

__global__ void hello_from_gpu()
{
const int blockid = blockIdx.x;
const int threadid = threadIdx.x;
printf("block index %d and thread idex %d!\n", blockid, threadid);
}
int kernel_apply1(void)
{
hello_from_gpu << <6, 5 >> > ();
cudaDeviceSynchronize();
return 0;
}

二、grid与block维度配置方法

使用kernel核函数可使用以下三种方法给出grid与block的维度配置。
代码如下：

　　 dim3 dimGrid(1);
    dim3 dimBlock(m);
    VecAdd1 << <dimGrid, dimBlock >> > (g_a, g_b, g_c); //方法一
    VecAdd1 << <dim3(1), dim3(m) >> > (g_a, g_b, g_c); //方法二
　　 VecAdd1 << <1, m >> > (g_a, g_b, g_c);  //方法三

三、基于1block与1third的应用

在grid中构建一个block，在block中构建一个third的方法demo。实际可理解为只有一个block，且该block只有一个x方向的维度，因此直接使用threadIdx.x处理，具体代码如下：

kernel核代码如下：

__global__ void VecAdd1(int* A, int* B, int* C)
{
    int i = threadIdx.x;
    C[i] = A[i] + B[i];
}

整体调用代码如下：

int kernel_apply2()
{
    int m = 8;    
    int* a, * b, *c;
    //分配host内存
    cudaMallocHost((void**)&a, sizeof(int) * m);
    cudaMallocHost((void**)&b, sizeof(int) * m);
    cudaMallocHost((void**)&c, sizeof(int) * m);

    std::cout << "value of a:" << endl;
    for (int i = 0; i < m; i++) {
        a[i] = rand() % 256;
        std::cout << a[i]<< "\t";
    }
    std::cout << "\nvalue of b:" << endl;
    for (int i = 0; i < m; i++) {
        b[i] = rand() % 260;
        std::cout << b[i]<< "\t";
    }

    int* g_a, * g_b,*g_c;
    //分配gpu内存
    cudaMalloc((void**)&g_a, sizeof(int) * m );
    cudaMalloc((void**)&g_b, sizeof(int) * m );
    cudaMalloc((void**)&g_c, sizeof(int) * m);
    // 赋值
    cudaMemcpy(g_a, a, sizeof(int) * m , cudaMemcpyHostToDevice);
    cudaMemcpy(g_b, b, sizeof(int) * m , cudaMemcpyHostToDevice);

    dim3 dimGrid(1);
    dim3 dimBlock(m);

    //应用grid只有x方向一个block，block只有x方向m个third
    VecAdd1 << <dimGrid, dimBlock >> > (g_a, g_b, g_c);
    //VecAdd1 << <dim3(1), dim3(m) >> > (g_a, g_b, g_c);
    //VecAdd1 << <1, m >> > (g_a, g_b, g_c);

    //将g_c赋值给c
    cudaMemcpy(c, g_c, sizeof(int) * m, cudaMemcpyDeviceToHost);
    //打印
    std::cout << "\nvalue of c:" << endl;
    for (int i = 0; i < m; i++) {
        std::cout <<c[i]<< "\t";
    }
    //释放内存
    cudaFree(g_a);
    cudaFree(g_b);
    cudaFree(g_c);
    cudaFreeHost(a);
    cudaFreeHost(b);
    cudaFreeHost(c);

    return 0;
}

运行结果如下：

四、基于1block与1third的深入探索

在grid中构建一个block，在block中构建一个third的方法demo的继续探索。kernel可以通过一维数组替换指针。

kernel核代码如下：

__global__ void MatAdd2(int A[8], int B[8], int C[8])
{     
    int i = threadIdx.x; 
    C[i] = A[i]+B[i];
}

整体代码：

int kernel_apply3()
{
    const int m = 8;
    int a[m],  b[m],  c[m];
    //int *a, *b, *c;
    //int* a, * b, c[m];

    //分配host内存
    cudaMallocHost((void**)&a, sizeof(int) * m);
    cudaMallocHost((void**)&b, sizeof(int) * m);
    cudaMallocHost((void**)&c, sizeof(int) * m);

    std::cout << "value of a:" << endl;
    for (int i = 0; i < m; i++) {
        a[i] = rand() % 69;
        std::cout << a[i] << "\t";
    }
        
    std::cout << "value of b:" << endl;
    for (int j = 0; j < m; j++) {
        b[j] = rand() % 25;
        std::cout << b[j] << "\t";
    }

    int *g_a, *g_b, *g_c;
    
    //分配gpu内存
    cudaMalloc((void**)&g_a, sizeof(int) * m);
    cudaMalloc((void**)&g_b, sizeof(int) * m);
    cudaMalloc((void**)&g_c, sizeof(int) * m);
    // 赋值
    cudaMemcpy(g_a, a, sizeof(int) * m, cudaMemcpyHostToDevice);
    cudaMemcpy(g_b, b, sizeof(int) * m, cudaMemcpyHostToDevice);

    MatAdd2 << <1, m >> > (g_a, g_b, g_c);
    cudaMemcpy(c, g_c, sizeof(int) * m, cudaMemcpyDeviceToHost);
    std::cout << "value of c:" << endl;
        for (int j = 0; j < m; j++) {  
            std::cout << c[j] << "\t";
        }
    return 0;
}

五、host端变量形式探索

方法一使用数组形式作为host端值；
方法二使用指针形式作为host端值，该方法较为常用；
方法三混合数组与指针方式作为host端值；

    int a[m],  b[m],  c[m];	方法一
    int *a, *b, *c;			方法二
    int* a, * b, c[m];		方法三

六、device端变量形式探索

我们以为device端也可类似host端一样有三种方式，然而使用以下方式，代码无法运行。

int a[m], b[m], c[m]; 
int* a, * b, c[m];

只能使用以下形式，方可使用。

int *g_a, *g_b, *g_c;

七、RGB图像转Gray图像

已介绍很多关于cuda应用，接下来，我们将介绍一个图像处理方法。

1、RGB转GRAY核函数编码方法

由于d_in可看成是2维数据，idx在cuda中表示为col，在图像中表示图像宽，idy表示逻辑类似idx，idy * imgwidth + idx可看做idy不变idx从0到imgwidth-1的取值。实际和上一章原理有些类似。

代码如下：

//用于CV读取图片BGR通道将其改为RGB方法
__global__ void rgb2grayincuda(uchar3* const d_in, unsigned char* const d_out,
    uint imgheight, uint imgwidth)
{
    const unsigned int idx = blockIdx.x * blockDim.x + threadIdx.x;  //w
    const unsigned int idy = blockIdx.y * blockDim.y + threadIdx.y;  //h

    if (idx < imgwidth && idy < imgheight)  //有的线程会跑到图像外面去，不执行即可
    {
        uchar3 rgb = d_in[idy * imgwidth + idx];
        d_out[idy * imgwidth + idx] = 0.299f * rgb.x + 0.587f * rgb.y + 0.114f * rgb.z;
   
    }
}

为了显示图像处理结果，使用以下代码显示查看图像处理结果，代码如下：

void show_img(Mat img) {
    cv::imshow("Image", img);
    cv::waitKey(1000);
    cv::destroyAllWindows();
}

2、cv::Mat复制device方法

grayImage为cv Mat矩阵，因此代码无需另外建立指针，而是通过grayImage.data可实现，因.data是uchar类型的指针，指向Mat数据矩阵的首地址，如下示列：

cudaMemcpy(grayImage.data, d_out, imgheight * imgwidth * sizeof(unsigned char), cudaMemcpyDeviceToHost);

以上代码是mat矩阵如何复制到gpu变量的方法，也是后面基于cpu版本的yolov5部署的cv mat数据转gpu的方法。

3、完整代码

oid kernel_apply4() {
    Mat srcImage = imread("./1.jpg");
    show_img(srcImage);
    const uint imgheight = srcImage.rows;
    const uint imgwidth = srcImage.cols;

    Mat grayImage(imgheight, imgwidth, CV_8UC1, Scalar(0));

    uchar3* d_in;   //向量类型，3个uchar
    unsigned char* d_out;

    cudaMalloc((void**)&d_in, imgheight * imgwidth * sizeof(uchar3));
    cudaMalloc((void**)&d_out, imgheight * imgwidth * sizeof(unsigned char));

    cudaMemcpy(d_in, srcImage.data, imgheight * imgwidth * sizeof(uchar3), cudaMemcpyHostToDevice);
    //说明：(imgwidth + threadsPerBlock.x - 1) / threadsPerBlock.x表示x方向
    dim3 threadsPerBlock(32, 32);
    dim3 blocksPerGrid((imgwidth + threadsPerBlock.x - 1) / threadsPerBlock.x, (imgheight + threadsPerBlock.y - 1) / threadsPerBlock.y);
    //启动内核
    rgb2grayincuda << <blocksPerGrid, threadsPerBlock >> > (d_in, d_out, imgheight, imgwidth);

    //执行内核是一个异步操作，因此需要同步以测量准确时间
    cudaDeviceSynchronize();
    //拷贝回来数据
    cudaMemcpy(grayImage.data, d_out, imgheight * imgwidth * sizeof(unsigned char), cudaMemcpyDeviceToHost);
    //释放显存
    cudaFree(d_in);
    cudaFree(d_out);

    imshow("grayImage", grayImage);
    cv::waitKey(1000);
    cv::destroyAllWindows();
}

八、总结

以上为探索kernel使用相关规则的一些补充，而本节重点应该是掌握如何使用kernel方法将图像转gray的应用，这也为后面最后基于cpu和gpu对yolov5的部署做一点衔接提示。