cu文件的编写
- .h文件的编写
- .cpp
- .cu
cuda编程手册地址, cuda基础
.h文件的编写
#ifndef TINYTRT_DECODE_H
#define TINYTRT_DECODE_H
#include "Trt.h"
#include "utils.h"
#include "YoloLayerPlugin/YoloLayerPlugin.hpp"
using namespace std;
struct Bbox {
int left, right, top, bottom;int clsId;float score;
};
cudaError_t decode_gpu(vector<float> input,YoloKernel yolo_kernel,vector<Detection>& output);
float my_decode(vector<float> intput,YoloKernel yolo_kernel,vector<Detection>& output);
#endif //TINYTRT_DECODE_H可以看出decode.h只有两个函数一个是cpu段的解码,一个是gpu段的解码,可以将cpu段解码的代码写进decode.cu文件中,也可以单独写进decode.cpp中.
.cpp
#include "decode.h"
float my_Logist(float data){
return 1./(1. + exp(-data)); }
float my_decode(vector<float> intput,YoloKernel yolo_kernel,vector<Detection>& output)
{
YoloKernel yolo=yolo_kernel;int stride=yolo.width*yolo.height;for(int i=0;i<yolo.width*yolo.height;i++){
for(int j=0;j<3;j++){
int begin_id=7*stride*j+i;int obj_id=begin_id+4*stride;float obj_prob=my_Logist(intput[obj_id]);if(obj_prob<0.7)continue;int class_id=-1;float max_prob=0.7;for(int k=0;k<2;k++){
float temp_prob=my_Logist(intput[begin_id+(5+k)*stride])*obj_prob;if(temp_prob>max_prob){
class_id=k;max_prob=temp_prob;}if(class_id>=0){
Detection det;int row=i/yolo.width;int cols=i%yolo.height;float a=my_Logist(intput[begin_id]);float b=my_Logist(intput[begin_id+stride]);det.bbox[0]=(cols+a)/yolo.width;det.bbox[1]=(row+b)/yolo.height;det.bbox[2]=exp(intput[begin_id+2*stride])*yolo.anchors[2*j];det.bbox[3]=exp(intput[begin_id+3*stride])*yolo.anchors[2*j+1];det.classId=class_id;det.prob=max_prob;output.emplace_back(det);}}}}
}
cpu端的解码没有什么可讲的,理解模型最后输出的是什么就可以.
模型最后的输出是一个batch*(channel*(box+1+class))yolo.wyolo.h的float数组,box指的是x,y,w,h.1代表的是这个anchor内有没有目标,class代表的是种类.
.cu
#include "decode.h"
#include "Trt.h"
#include "utils.h"
#include "math.h"
#include <iostream>
#include <fstream>
#include <cmath>
#include "../plugin/YoloLayerPlugin/YoloLayerPlugin.hpp"
__device__ float Logist1(float data){
return 1./(1. + exp(-data)); };
__global__ void caldetection(const float* input,float* output,int noelements,int yolowidth,int yoloheight,const float anchors[6],int classes,int outputElem)
{
int idx=threadIdx.x+blockDim.x*blockIdx.x;if(idx>noelements)return;int stride=yoloheight*yolowidth;int bnidx=idx/stride;int curidx=idx-bnidx*stride;const float* curinput=input+bnidx*(7)*stride*3;for(int k=0;k<3;k++){
int beginidx=(7*stride)*k+curidx;int objidx=beginidx+stride*4;float objprob=Logist1(curinput[objidx]);if(objprob <= 0.7)continue;int row = curidx / yolowidth;int cols = curidx % yolowidth;int classId = -1;float maxProb = IGNORE_THRESH;for (int c = 0;c<2;++c){
float cProb = Logist1(curinput[beginidx + (5 + c) * stride]) * objprob;if(cProb > maxProb){
maxProb = cProb;classId = c;}}if(classId >= 0) {
float *curOutput = output + bnidx*outputElem;int resCount = (int)atomicAdd(curOutput,1);char* data = (char * )curOutput + sizeof(float) + resCount*sizeof(Detection);Detection* det = (Detection*)(data);det->bbox[0] = (cols + Logist1(curinput[beginidx]))/ yolowidth;det->bbox[1] = (row + Logist1(curinput[beginidx+stride]))/ yoloheight;det->bbox[2] = exp(curinput[beginidx+2*stride]) * anchors[2*k];det->bbox[3] = exp(curinput[beginidx+3*stride]) * anchors[2*k + 1];float tem_cla=float(classId);det->classId = llround(double(tem_cla));det->prob = maxProb;}}}
cudaError_t decode_gpu(vector<float> input,YoloKernel yolo_kernel,vector<Detection>& output)
{
float* temp_input;int input_num=input.size();CUDA_CHECK(cudaMalloc(&temp_input,input.size()*sizeof(float)));CUDA_CHECK(cudaMemcpy(temp_input,&input[0],input.size()*sizeof(float),cudaMemcpyHostToDevice));float* output1;void* devAnchor;size_t AnchorLen = sizeof(float)* CHECK_COUNT*2;CUDA_CHECK(cudaMalloc(&devAnchor,AnchorLen));int outputElem = 1;outputElem+=yolo_kernel.width*yolo_kernel.height*3*sizeof(Detection)/sizeof(float);CUDA_CHECK(cudaMalloc(&output1,sizeof(float)*outputElem));int numelem=yolo_kernel.width*yolo_kernel.height;CUDA_CHECK(cudaMemcpyAsync(devAnchor,yolo_kernel.anchors,AnchorLen,cudaMemcpyHostToDevice));caldetection<<<(yolo_kernel.width*yolo_kernel.height+512-1)/512,512>>>(temp_input,output1,numelem,yolo_kernel.width,yolo_kernel.height,(float *)devAnchor,2,outputElem);cudaError_t cudaStatus;cudaFree(devAnchor);float* out_host{
};CUDA_CHECK(cudaMallocHost(&out_host,sizeof(float)*outputElem));CUDA_CHECK(cudaMemcpy(out_host,output1,sizeof(float)*outputElem,cudaMemcpyDeviceToHost));
// printf("第一个输出%f,第二个输出%f",out_host[0],out_host[1]);for(int k=0;k<int(out_host[0]);k++){
Detection temp;temp.bbox[0]=out_host[6*k+1];temp.bbox[1]=out_host[6*k+2];temp.bbox[2]=out_host[6*k+3];temp.bbox[3]=out_host[6*k+4];temp.classId=out_host[6*k+5];temp.prob=out_host[6*k+6];output.push_back(temp);}cudaStatus = cudaGetLastError();if (cudaStatus != cudaSuccess) {
std::cerr << "CUDA error " << cudaGetErrorString(cudaStatus) << " at " << __FILE__ << ":" << __LINE__ << std::endl;}return cudaGetLastError();.cu文件中要注意,kernel函数中不能debug进去,也不能将中间结果输出到txt中查看,(目前水平是这样),因此只能靠printf调试,还是很麻烦的.
注意__device__,global,__host__的区别,__device__由gpu调用在gpu上执行,global__由cpu调用在GPU上执行.所以decode.cu的核函数为__global,再进去__global__函数前,要将cpu中的数据拷贝到gpu,还要在gpu上申请输出的空间块(这个空间要大于等于你所需要的空间)
这一块指针,空间申请之类的c++,cuda语法还要多多研究,先谈点现在简单的理解.
float* temp_input;CUDA_CHECK(cudaMalloc(&temp_input,input.size()*sizeof(float)));CUDA_CHECK(cudaMemcpy(temp_input,&input[0],input.size()*sizeof(float),cudaMemcpyHostToDevice));void* devAnchor;size_t AnchorLen = sizeof(float)* CHECK_COUNT*2;CUDA_CHECK(cudaMalloc(&devAnchor,AnchorLen));CUDA_CHECK(cudaMemcpy(devAnchor,yolo_kernel.anchors,AnchorLen,cudaMemcpyHostToDevice));float* output1;int outputElem = 1;outputElem+=yolo_kernel.width*yolo_kernel.height*3*sizeof(Detection)/sizeof(float);CUDA_CHECK(cudaMalloc(&output1,sizeof(float)*outputElem));
关于cudaMalloc()和cudaMemcpy()函数的使用详见cuda编程手册.
注意核函数中要用到的所有host数据都要cpy到device在送入核函数.
float* out_host{
};
CUDA_CHECK(cudaMallocHost(&out_host,sizeof(float)*outputElem));
CUDA_CHECK(cudaMemcpy(out_host,output1,sizeof(float)*outputElem,cudaMemcpyDeviceToHost));
这一步是将device中的值cpy到host处理,要不gpu中的值是不可见状态.
caldetection<<<(yolo_kernel.width*yolo_kernel.height+512-1)/512,512>>>(temp_input,output1,numelem,yolo_kernel.width,yolo_kernel.height,(float *)devAnchor,2,outputElem);
注意下kernel函数的调用,使用的是<<<block,thread>>>形式,其block,thread可以是一维,也可以是二维,每个的idx是确认好的.核函数简单理解就是同时进行很多for循环,这一部分还要多看多写.
cudaFree(devAnchor)
cudaFree(output1);
cudaFreeHost(out_host);
最后使用完要记得释放指针,避免报段错误.
int nByte=sizeof(float)*nElem;
float *res_h=(float*)malloc(nByte);
memset(res_h,0,nByte);
free(res_h);
常见的c++内存拷贝
float *curOutput = output + bnidx*outputElem;
int resCount = (int)atomicAdd(curOutput,1);
char* data = (char * )curOutput + sizeof(float) + resCount*sizeof(Detection);
Detection* det = (Detection*)(data);
det->bbox[0] = (cols + Logist1(curinput[beginidx]))/ yolowidth;
注意下atomicADD()的使用,大概是用来多线程计数的.指针数组申请时一定要分配空间块大小及地址.还有强制类型转换的用法.