diff --git a/CMakeLists.txt b/CMakeLists.txt index a8f4744..225c9e1 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -31,7 +31,8 @@ cuda_add_library(kernels SHARED src/kernels/activation_elu.cu src/kernels/activation_leaky.cu src/kernels/activation_logistic.cu src/kernels/reorg.cu - src/kernels/softmax.cu) + src/kernels/softmax.cu + src/kernels/convert.cu) file(GLOB tkdnn_SRC "src/*.cpp") set(tkdnn_LIBS kernels ${CUDA_LIBRARIES} ${CUDA_CUBLAS_LIBRARIES} -lcudnn -lnvinfer ${OpenCV_LIBS}) diff --git a/demo/detection/detection.cpp b/demo/detection/detection.cpp index f0b78c7..838cf88 100644 --- a/demo/detection/detection.cpp +++ b/demo/detection/detection.cpp @@ -46,7 +46,8 @@ cv::Mat GetSquareImage(const cv::Mat& img, int target_width) { return square; } -void compute_image( cv::Mat imageORIG, +//return inference time +double compute_image( cv::Mat imageORIG, tkDNN::NetworkRT *netRT, tkDNN::RegionInterpret *rI, dnnType *input, dnnType *output) { @@ -83,11 +84,14 @@ void compute_image( cv::Mat imageORIG, rI->interpretData(output, imageORIG.cols, imageORIG.rows); + return t_ns; } int print_usage() { - std::cout<<"usage: ./detection net.rt validation_list.txt [-t ] [-s]\n" - <<" -t: set thresh value\n -s: show images as compute\n\n" + std::cout<<"usage: ./detection net.rt validation_list.txt" + <<" [-t ] [-s] [-i ]\n" + <<" -t: set thresh value\n -s: show images as compute\n" + <<" -i: images to compute\n\n" <<"> validation_list.txt format: \n" <<" path/to/image.jpg path/to/label.txt\n" <<"> label.txt format: \n" @@ -105,13 +109,15 @@ int main(int argc, char *argv[]) { char *imageset_path = NULL; float thresh = 0.3f; bool show = false; + int iterations = INT_MAX; //parse params int c; - while ((c = getopt (argc, argv, "t:s")) != -1) { + while ((c = getopt (argc, argv, "t:si:")) != -1) { switch(c) { - case 't': thresh = atof(optarg); break; - case 's': show = true; break; + case 't': thresh = atof(optarg); break; + case 's': show = true; break; + case 'i': iterations = atoi(optarg); break; case '?': return print_usage(); default: return print_usage(); @@ -141,9 +147,13 @@ int main(int argc, char *argv[]) { if(!imageset.is_open()) FatalError("could not read imageset"); + double mTime = 0; float mAP = 0; int processed_images; - for(processed_images=1; getline(imageset, line); processed_images++) { + + for(processed_images=1; + processed_images-1 < iterations && getline(imageset, line); + processed_images++) { std::string image_path = line.substr(0, line.find(" ")); std::string label_path = line.substr(line.find(" ")+1, line.size()); @@ -155,7 +165,7 @@ int main(int argc, char *argv[]) { FatalError("Could not open image"); std::cout<<"Image size: ("< +#include #include "Layer.h" +#include "kernels.h" namespace tkDNN { @@ -26,6 +28,72 @@ LayerWgs::LayerWgs(Network *net, int inputs, int outputs, readBinaryFile(weights_path.c_str(), outputs, &mean_h, &mean_d, seek); seek += outputs; readBinaryFile(weights_path.c_str(), outputs, &variance_h, &variance_d, seek); + + float eps = CUDNN_BN_MIN_EPSILON; + + power_h = new dnnType[outputs]; + for(int i=0; ifp16) + return; + + //convert to fp16 + int w_size = inputs*outputs*kh*kw*kl; + data16_h = new __half[w_size]; + cudaMalloc(&data16_d, w_size*sizeof(__half)); + float2half(data_d, data16_d, w_size); + cudaMemcpy(data16_h, data16_d, w_size*sizeof(__half), cudaMemcpyDeviceToHost); + + int b_size = outputs; + bias16_h = new __half[b_size]; + cudaMalloc(&bias16_d, w_size*sizeof(__half)); + float2half(bias_d, bias16_d, b_size); + cudaMemcpy(bias16_h, bias16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost); + + if(batchnorm) { + + power16_h = new __half[b_size]; + mean16_h = new __half[b_size]; + variance16_h = new __half[b_size]; + scales16_h = new __half[b_size]; + + cudaMalloc(&power16_d, b_size*sizeof(__half)); + cudaMalloc(&mean16_d, b_size*sizeof(__half)); + cudaMalloc(&variance16_d, b_size*sizeof(__half)); + cudaMalloc(&scales16_d, b_size*sizeof(__half)); + + //temporary buffers + float *tmp_d; + cudaMalloc(&tmp_d, b_size*sizeof(float)); + + //init power array of ones + cudaMemcpy(tmp_d, power_h, b_size*sizeof(float), cudaMemcpyHostToDevice); + float2half(tmp_d, power16_d, b_size); + cudaMemcpy(power16_h, power16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost); + + //mean array + + cudaMemcpy(tmp_d, mean_h, b_size*sizeof(float), cudaMemcpyHostToDevice); + float2half(tmp_d, mean16_d, b_size); + cudaMemcpy(mean16_h, mean16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost); + + //convert variance + + cudaMemcpy(tmp_d, variance_h, b_size*sizeof(float), cudaMemcpyHostToDevice); + float2half(tmp_d, variance16_d, b_size); + cudaMemcpy(variance16_h, variance16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost); + + //conver scales + float2half(scales_d, scales16_d, b_size); + cudaMemcpy(scales16_h, scales16_d, b_size*sizeof(__half), cudaMemcpyDeviceToHost); } } diff --git a/src/Network.cpp b/src/Network.cpp index 7157293..3cc342d 100644 --- a/src/Network.cpp +++ b/src/Network.cpp @@ -1,5 +1,5 @@ #include -#include +#include #include "tkdnn.h" #include "Network.h" @@ -22,6 +22,14 @@ Network::Network(dataDim_t input_dim) { checkERROR( cublasCreate(&cublasHandle) ); num_layers = 0; + + fp16 = false; + if(const char* env_p = std::getenv("TKDNN_MODE")) + if(strcmp(env_p, "FP16") == 0) + fp16 = true; + + if(fp16) + std::cout< #include #include +#include // memcpy +#include +#include "kernels.h" + +#include "utils.h" #include "NvInfer.h" - #include "NetworkRT.h" using namespace nvinfer1; @@ -45,7 +49,7 @@ NetworkRT::NetworkRT(Network *net, const char *name) { builderRT->setMaxBatchSize(1); builderRT->setMaxWorkspaceSize(1 << 30); -/* + //change datatype based on system specs if(builderRT->platformHasFastInt8()) { BatchStream bstream({32,dim.c, dim.h, dim.w}, 32, 1); @@ -53,13 +57,13 @@ NetworkRT::NetworkRT(Network *net, const char *name) { builderRT->setInt8Mode(true); builderRT->setInt8Calibrator(&calib); - } else if(builderRT->platformHasFastFp16()) { + } else if(net->fp16 && builderRT->platformHasFastFp16()) { dtRT = DataType::kHALF; builderRT->setHalf2Mode(true); } -*/ + //add input layer - ITensor *input = networkRT->addInput("data", dtRT, + ITensor *input = networkRT->addInput("data", DataType::kFLOAT, DimsCHW{ dim.c, dim.h, dim.w}); checkNULL(input); @@ -170,22 +174,49 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Layer *l) { ILayer* NetworkRT::convert_layer(ITensor *input, Dense *l) { //std::cout<<"convert Dense\n"; + void *data_b, *bias_b; + if(dtRT == DataType::kHALF) { + data_b = l->data16_h; + bias_b = l->bias16_h; + } else { + data_b = l->data_h; + bias_b = l->bias_h; + } - Weights w { dtRT, l->data_h, l->inputs*l->outputs}; - Weights b = { dtRT, l->bias_h, l->outputs}; + Weights w { dtRT, data_b, l->inputs*l->outputs}; + Weights b = { dtRT, bias_b, l->outputs}; IFullyConnectedLayer *lRT = networkRT->addFullyConnected(*input, l->outputs, w, b); checkNULL(lRT); return lRT; } + ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { //std::cout<<"convert conv2D\n"; - Weights w { dtRT, l->data_h, l->inputs*l->outputs*l->kernelH*l->kernelW}; + void *data_b, *bias_b, *power_b, *mean_b, *variance_b, *scales_b; + if(dtRT == DataType::kHALF) { + data_b = l->data16_h; + bias_b = l->bias16_h; + power_b = l->power16_h; + mean_b = l->mean16_h; + variance_b = l->variance16_h; + scales_b = l->scales16_h; + } else { + data_b = l->data_h; + bias_b = l->bias_h; + power_b = l->power_h; + mean_b = l->mean_h; + variance_b = l->variance_h; + scales_b = l->scales_h; + } + + + Weights w { dtRT, data_b, l->inputs*l->outputs*l->kernelH*l->kernelW}; Weights b; if(!l->batchnorm) - b = { dtRT, l->bias_h, l->outputs}; + b = { dtRT, bias_b, l->outputs}; else b = { dtRT, nullptr, 0}; //on batchnorm bias are added later @@ -198,29 +229,15 @@ ILayer* NetworkRT::convert_layer(ITensor *input, Conv2d *l) { lRT->setPadding(DimsHW{l->paddingH, l->paddingW}); if(l->batchnorm) { - float eps = CUDNN_BN_MIN_EPSILON; - - //make power array of ones - dnnType *power_h = new dnnType[l->outputs]; - for(int i=0; ioutputs; i++) power_h[i] = 1.0f; - - //convert mean - for(int i=0; ioutputs; i++) - l->mean_h[i] = l->mean_h[i] / -sqrt(eps + l->variance_h[i]); - - //convert variance - for(int i=0; ioutputs; i++) - l->variance_h[i] = 1.0f / sqrt(eps + l->variance_h[i]); - - Weights power{dtRT, power_h, l->outputs}; - Weights shift{dtRT, l->mean_h, l->outputs}; - Weights scale{dtRT, l->variance_h, l->outputs}; + Weights power{dtRT, power_b, l->outputs}; + Weights shift{dtRT, mean_b, l->outputs}; + Weights scale{dtRT, variance_b, l->outputs}; IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, shift, scale, power); checkNULL(lRT2); - Weights shift2{dtRT, l->bias_h, l->outputs}; - Weights scale2{dtRT, l->scales_h, l->outputs}; + Weights shift2{dtRT, bias_b, l->outputs}; + Weights scale2{dtRT, scales_b, l->outputs}; IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, shift2, scale2, power); checkNULL(lRT3); diff --git a/src/kernels/convert.cu b/src/kernels/convert.cu new file mode 100644 index 0000000..f7a7f72 --- /dev/null +++ b/src/kernels/convert.cu @@ -0,0 +1,21 @@ +#include "kernels.h" + +__global__ +void float2half_device(float *input, __half *output, int size) { + + int i = blockDim.x*blockIdx.x + threadIdx.x; + + if(i>>(srcData, dstData, size); + cudaDeviceSynchronize(); +} diff --git a/src/utils.cpp b/src/utils.cpp index 13c93b1..f250238 100644 --- a/src/utils.cpp +++ b/src/utils.cpp @@ -70,7 +70,7 @@ void printDeviceVector(int size, dnnType* vec_d, bool device) int checkResult(int size, dnnType *data_d, dnnType *correct_d, bool device) { dnnType *data_h, *correct_h; - const float eps = 0.0001f; + const float eps = 0.001f; if(device) { data_h = new dnnType[size]; diff --git a/tests/yolo_relu/yolo_relu.cpp b/tests/yolo_relu/yolo_relu.cpp index ecc4862..be1294a 100644 --- a/tests/yolo_relu/yolo_relu.cpp +++ b/tests/yolo_relu/yolo_relu.cpp @@ -139,12 +139,13 @@ int main() { std::cout<<"CUDNN vs correct"; checkResult(out_dim, out_data, out); std::cout<<"TRT vs correct"; checkResult(out_dim, out_data2, out); std::cout<<"CUDNN vs TRT "; checkResult(out_dim, out_data, out_data2); - + std::cout<<"\n\nDetected objects: \n"; dnnType *output_h = new dnnType[rI.output_dim.tot()]; checkCuda(cudaMemcpy(output_h, out_data2, rI.output_dim.tot()*sizeof(dnnType), cudaMemcpyDeviceToHost)); rI.interpretData(output_h, 608, 608); rI.showImageResult(input_h); + return 0; }