cuda系列详细教程

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

链接：https://blog.csdn.net/weixin_38252409/category_12383040.html?spm=1001.2014.3001.5482

 一、核函数index寻找

cuda通过线程执行并行运算，理所当然，我们需要知道如何使用每个线程实现自己的计算逻辑。而线程操作通过索引(index)操作，索引和block与grid挂钩，自然我们需要知晓如何在grid与block中确定索引。为此，我写了索引寻找规律，可以通过公式直接计算，我在此不细说，仅展示以下展示部分代码，其详细内容和附件可点击我的链接。

部分展示代码如下：

3d grid与1d block寻找索引代码：

blockSize = blockDim.x（一维 block 的大小）
blockId = Dx * Dy * z + Dx * y + x （三维 grid 中 block 的 id，用公式）
　　= gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x
threadId = threadIdx.x （一维 block 中 thread 的 id）
Id = (gridDim.x * gridDim.y * blockIdx.z + gridDim.x * blockIdx.y + blockIdx.x ) * blockDim.x + threadIdx.x

1d grid, 2d block寻找索引代码：

blockSize = blockDim.x * blockDim.y（二维 block 的大小）
blockId = blockIdx.x（一维 grid 中 block id）
threadId = Dx * y + x （二维 block 中 thread 的 id）
　　= blockDim.x * threadIdx.y + threadIdx.x
Id = blockIdx.x * (blockDim.x * blockDim.y) + blockDim.x * threadIdx.y + threadIdx.x

二、kernel函数实例

　　如上所说，我们知道kernel函数索引寻找方法，我们自然想通过索引实现各种运算，多数为矩阵运算。为此，我写了大量实例cuda代码，并用不同实例说明其cuda编码规律，我将以其中一个实例矩阵加法代码作为展示，此代码使用多种途径求其结果，其详细内容和附件代码可点击我的链接。

__global__ void gpu_matrix_plus_thread(int* a, int* b, int* c)
{   
    //方法一：通过id方式计算
    //grid为2维度，block为2维度,使用公式id=blocksize * blockid + threadid
    int blocksize = blockDim.x*blockDim.y;
    int blockid = gridDim.x*blockIdx.y+blockIdx.x;
    int threadid = blockDim.x*threadIdx.y+threadIdx.x;
    int id = blocksize * blockid + threadid;    
    c[id] = a[id] + b[id];
}
__global__ void gpu_matrix_plus1(int* a, int* b, int* c, int m, int n)
{   //方法二：通过row与col的方式计算-->通过变换列给出id
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    c[row*n  + col] = a[row*n  + col] + b[row*n  + col];
}
__global__ void gpu_matrix_plus2(int* a, int* b, int* c, int m, int n)
{   //方法三：通过row与col的方式计算-->通过变换行给出id
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    c[row  + col*m] = a[row  + col*m] + b[row  + col*m];
}

三、性能优化(内存)

　　既然我们已可自如实现自己计算逻辑，那么我么也需兼顾运行效率，运行速度提升可通过更好pipeline逻辑实现，也可通过内存方式实现，而逻辑架构因人而异，我将不在细说，内存实现可通过对cuda内存理解便可掌握。为此，我也写了内存相关实例代码，介绍其内存使用方法，如纹理内存、共享内存等，我将以纹理内存代码作为展示，其更多详细内容和附件代码可点击我的链接。

//核心代码，在gpu端执行的kernel，  
__global__ void Textureone(unsigned int* listTarget, int size)
{
    unsigned int texvalue = 0;
    int index = blockIdx.x * blockDim.x + threadIdx.x; //通过线程ID得到数组下标 
    if (index < size)
        texvalue= tex1Dfetch(texone, index)*100; //通过索引获得纹理值再乘100 
        listTarget[index] = texvalue;
}

四、原子操作

　　在 CUDA 中，原子操作是一种用于确保多个线程同时访问同一内存地址时的同步机制。原子操作可以确保只有一个线程可以访问内存地址，并且可以避免数据竞争和不确定的结果。我将已实例代码展示原子操作，以下为原子操作实例，其更多详细内容和附件代码可点击我的链接。

部分原子操作代码如下：

__global__ void kernel(int* data) {
    int tid = threadIdx.x + blockIdx.x * blockDim.x;
    // 对共享内存中的数据执行原子加操作
    atomicAdd(&data[tid], 1);
}

int main() {
    int size = 1024;
    int* data = new int[size];
    int* d_data;
    cudaMalloc(&d_data, size * sizeof(int));
    cudaMemcpy(d_data, data, size * sizeof(int), cudaMemcpyHostToDevice);
    kernel<<<1, size>>>(d_data);
    cudaMemcpy(data, d_data, size * sizeof(int), cudaMemcpyDeviceToHost);
    cudaFree(d_data);
    delete[] data;
    return 0;
}

五、流stream

六、cuda处理nms编码

大量算法的后处理逻辑均会使用NMS算法去重，然CPU算法较慢，为此我写了cuda的NMS算法处理，以下将部分展示，其详细代码可点击我的链接。

部分代码如下：

// 定义CUDA核函数，用于执行NMS算法
__global__ void nms_kernel(nms_box* boxes, int* indices, int* num_indices, float nms_thr)
{
    /*
    boxes:输入nms信息，为结构体
    indices:输入为列表序列，记录所有box，如[0,1,2,3,4,5,...]，后续将不需要会变成-1。
    num_indices:记录有多少个box数量
    float nms_thr:nms的阈值，实际为iou阈值
    */
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= *num_indices) { return; }
    int index = indices[i];
    if (index == -1) { return; }
    nms_box box = boxes[index];
    for (int j = i + 1; j < *num_indices; j++) {
        int other_index = indices[j];
        if (other_index == -1) { continue; }
        nms_box other_box = boxes[other_index];
        float iou_value = iou(box, other_box);
        printf("iou value:%f\n", iou_value);
        if (iou_value > nms_thr) { indices[j] = -1; }
    }
}

输出结果如下：

七、cuda处理yolo算法输出编码

在tensorrt部署中，yolo算法输出在gpu设备上且数据较为庞大(如：640输入将有25200*(class_num+5)*batch数据)，使用cpu处理，需将值从gpu端复制host端，复制过程会花费很多时间。因此，我为此我写了cuda的yolo输出数据处理，以下将部分展示，其详细代码可点击我的链接。

代码如下：

__global__ void decode_yolo_kernel(float* prob, float* parray, int max_objects, int cls_num, float conf_thr, int* d_count) {
    int idx = blockDim.x * blockIdx.x + threadIdx.x; 
    int tmp_idx = idx * (cls_num + 5); 
    float left = prob[tmp_idx + 0]; 
    float top = prob[tmp_idx + 1];
    float right = prob[tmp_idx + 2];
    float bottom = prob[tmp_idx + 3];
    float conf = prob[tmp_idx + 4]; 
    float class_score = prob[tmp_idx + 5];
    float tmp_conf = conf * class_score;
    int class_id = 0;
    for (int j = 0; j < cls_num; j++) {
        int cls_idx = tmp_idx + 5 + j;
        if (tmp_conf < conf * prob[cls_idx]) {
            class_id = j;
            tmp_conf = conf * prob[cls_idx];
        }
    }
    if (tmp_conf < conf_thr) { return; }
    int index = atomicAdd(d_count, 1);
    if (index >= max_objects) { return; } 
    int out_index = index * 6; 
    parray[out_index + 0] = left;   
    parray[out_index + 1] = top;    
    parray[out_index + 2] = right;  
    parray[out_index + 3] = bottom;         parray[out_index + 4] = tmp_conf;   
    parray[out_index + 5] = class_id;   

}

代码解释：

目的：简化模拟yolo输出结果于cuda核函数中处理

假设：置信度阀值为0.45，类别为2，最大目标数为3，

下图数据说明：左边为核函数实现代码展示；右边每一含为一个目标预测结果，分别表示box值[x,y,w,h]、置信度conf、类别预测值[c1_score,c2_score]；右下角为cuda核函数选择目标结果，其中conf为类别score*conf；

实现方法：利用核函数与原子操作完成目标筛选。

运行结果如下：

八、cuda处理yolo算法整个过程

在tensorrt部署中，yolo算法输出使用gpu处理已在上面涉及，然如何去重box与yolo整套后处理呢？为此，我也写了基于cuda处理yolo的整个过程，以下将部分代码展示，其详细代码可点击我的链接。

代码如下：

void imitate_yolo_postprocess_convert() {
   
    const int block = 32;
    /*************************************************开始cuda计算***********************************************/
    cudaStream_t stream;
    cudaStreamCreate(&stream);

    h_count = 0;
    cudaMemcpy(d_count, &h_count, sizeof(int), cudaMemcpyHostToDevice);  //初始化记录有效变量d_count与h_count
    int grid = (anchor_output_num + block - 1) / block;
    decode_yolo_kernel << < grid, block, 0, stream >> > (gpu_input, gpu_output, max_object, cls_num, conf_thr, d_count);
    cudaMemcpy(&h_count, d_count, sizeof(int), cudaMemcpyDeviceToHost);
    if (h_count > max_object) { h_count = max_object; };


    /****************************************打印模型输出输出数据结果--》通过置信度已过滤不满足要求和给出类别**********************************/

    float* host_decode = nullptr; // 保存gpu处理的变量 
    cudaMallocHost((void**)&host_decode, sizeof(float) * max_object * 6);
    cudaMemcpy(host_decode, gpu_output, sizeof(float) * max_object * 6, cudaMemcpyDeviceToHost);
    std::cout << "\n\n打印输出结果-gpu_output\n" << endl;
    if (h_count == 0) { std::cout << "\n无检测结果" << endl; }
    for (int i = 0; i < h_count; i++) {
        int idx = i * 6;
        std::cout << "x1:" << host_decode[idx] << "\ty1:" << host_decode[idx + 1] << "\tx2:" << host_decode[idx + 2]
            << "\ty2:" << host_decode[idx + 3] << "\tconf:" << host_decode[idx + 4] << "\tclass_id:" << host_decode[idx + 5] << endl;

    }
    /******************************************************************************************************************************/

    int grid_max = (max_object + block - 1) / block;
    data_format_convert << < grid_max, block, 0, stream >> > (d_boxes, gpu_output, h_count); // gpu_output格式为[x1,y1,conf,cls_id]

    /****************************************将数据转换为带有nms_box格式数据******************************************************/
    nms_box* h_boxes_format = nullptr;
    cudaMallocHost(&h_boxes_format, anchor_output_num * sizeof(nms_box));
    cudaMemcpy(h_boxes_format, d_boxes, anchor_output_num * sizeof(nms_box), cudaMemcpyDeviceToHost);
    std::cout << "\n\n打印格式转换输出-h_boxes_format\n" << endl;
    if (h_count == 0) { std::cout << "\n无检测结果" << endl; }
    for (int i = 0; i < h_count; i++) {
        nms_box bb = h_boxes_format[i];
        std::cout << "x1:" << bb.x1 << "\ty1:" << bb.y1 << "\tx2:" << bb.x2 << "\ty2:" << bb.y2 << "\tconf:" << bb.score << "\tclass_id:" << bb.cls_id << endl;
    }
    /******************************************************************************************************************************/


    cudaMemcpy(d_nms_indices, h_nms_indices_init, max_object * sizeof(int), cudaMemcpyHostToDevice); //初始化nms处理的索引-->很重要

    /****************************************查看d_nms_indices数据******************************************************/
    int* d_nms_indices_visual = nullptr;
    cudaMallocHost(&d_nms_indices_visual, max_object * sizeof(int));
    cudaMemcpy(d_nms_indices_visual, d_nms_indices, max_object * sizeof(int), cudaMemcpyDeviceToHost);
    std::cout << "\n\nd_nms_indices:\n" << endl;
    for (int i = 0; i < max_object; i++) { std::cout << "\t" << d_nms_indices_visual[i] << endl; }

    /******************************************************************************************************************************/

    nms_yolo_kernel << <grid_max, block >> > (d_boxes, d_nms_indices, h_count, nms_thr);

    /*******将yolo的gpu上结果转host端，然后保存结果处理-->最终结果保存在keep_boxes中**********/
    cudaMemcpy(h_boxes, d_boxes, anchor_output_num * sizeof(nms_box), cudaMemcpyDeviceToHost);
    cudaMemcpy(h_nms_indices, d_nms_indices, max_object * sizeof(int), cudaMemcpyDeviceToHost);  //保存处理后的indice

    vector<nms_box> keep_boxes(h_count);
    for (int i = 0; i < h_count; i++) {
        if (h_nms_indices[i] > -1) {
            keep_boxes[i] = h_boxes[i];
        }
    }


    /****************************************查看nms处理后的-d_nms_indices******************************************************/
    std::cout << "nms处理后，保留box索引，-1表示排除obj，>-1表示保存obj" << endl;
    for (int i = 0; i < max_object; i++) { std::cout << h_nms_indices[i] << "\t"; }
    /**********************************************************************************************/

    /****************************************随便一张图为背景-显示结果于图上******************************************************/
    cv::Mat image = cv::imread("image.jpg");

    for (nms_box box : keep_boxes) {

        cv::Point p1(box.x1, box.y1);
        cv::Point p2(box.x2, box.y2);
        cv::rectangle(image, p1, p2, cv::Scalar(0, 255, 0), 4, 1, 0);//矩形的两个顶点，两个顶点都包括在矩形内部
    }


    cv::resize(image, image, cv::Size(600, 400), 0, 0, cv::INTER_NEAREST);
    cv::imshow("www", image);
    cv::waitKey(100000);
    cv::destroyAllWindows();
    /**********************************************************************************************/


}

注：以上代码有删减，完整代码可点击链接

结果显示如下：

九、yolo的tensorrt部署(前后处理的cpu版与gpu版)

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

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

链接：https://blog.csdn.net/weixin_38252409/category_12383040.html?spm=1001.2014.3001.5482