fp16 implementation, TODO deallocate in LayerWgs

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

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 <thresh>] [-s]\n"
+    std::cout<<"usage: ./detection net.rt validation_list.txt"
-             <<" -t: set thresh value\n -s: show images as compute\n\n" 
+             <<" [-t <thresh>] [-s] [-i <iterations>]\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 't':   thresh = atof(optarg);       break;
-        case 's':   show = true;           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: ("<<img.cols<<"x"<<img.rows<<")\n";
 
-        compute_image(img, &netRT, &rI, input, output); 
+        mTime += compute_image(img, &netRT, &rI, input, output); 
 
         std::ifstream labels(label_path.c_str());
         if(!labels.is_open())
@@ -223,9 +233,17 @@ int main(int argc, char *argv[]) {
         if(show) {
             cv::namedWindow("result");
             cv::imshow("result", img);
-            cv::waitKey(1000);
+            cv::waitKey(10);
         }        
     }
 
+    //print results to file
+    processed_images -= 1;
+    std::ofstream res("results.txt", std::ios::app);
+    res<<"#### "<<tensor_path<<"\n";
+    res<<"processed images: "<<processed_images<<"\n";
+    res<<"mean inference time: "<<mTime/processed_images<<"\n";
+    res<<"mean AP: "<<mAP/processed_images<<"\n";
+    res<<"thesh used: "<<thresh<<"\n\n";
     return 0;
 }

include/Layer.h

@@ -80,9 +80,19 @@ public:
 
     //batchnorm
     bool batchnorm;
+    dnnType *power_h;
     dnnType *scales_h,   *scales_d;
     dnnType *mean_h,     *mean_d;
     dnnType *variance_h, *variance_d;
+
+    //fp16
+    __half *data16_h, *bias16_h;
+    __half *data16_d, *bias16_d;
+
+    __half *power16_h,    *power16_d;
+    __half *scales16_h,   *scales16_d;
+    __half *mean16_h,     *mean16_d;
+    __half *variance16_h, *variance16_d;
 };
 
 

include/Network.h

@@ -58,7 +58,9 @@ public:
 
     dataDim_t input_dim;
     dataDim_t getOutputDim();
+
+    bool fp16;
 };
 
 }
-#endif //NETWORK_H
+#endif //NETWORK_H

include/kernels.h

@@ -12,4 +12,6 @@ void reorgForward(  dnnType* srcData, dnnType* dstData,
 void softmaxForward(float *input, int n, int batch, int batch_offset, 
                     int groups, int group_offset, int stride, float temp, float *output, cudaStream_t stream = cudaStream_t(0));
 
+
+void float2half(float* srcData, __half* dstData, int size, const cudaStream_t stream = cudaStream_t(0));
 #endif //KERNELS_H

src/LayerWgs.cpp

@@ -1,6 +1,8 @@
 #include <iostream>
+#include <string.h>
 
 #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; i<outputs; i++) power_h[i] = 1.0f;
+
+        for(int i=0; i<outputs; i++)
+            mean_h[i] = mean_h[i] / -sqrt(eps + variance_h[i]); 
+
+        for(int i=0; i<outputs; i++)
+            variance_h[i] = 1.0f / sqrt(eps + variance_h[i]);
+    }
+
+
+    if(!net->fp16)
+        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);
     }
 }
 

src/Network.cpp

@@ -1,5 +1,5 @@
 #include <iostream>
-#include <string>
+#include <string.h>
 
 #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<<COL_REDB<<"!! FP16 INERENCE ENABLED !!"<<COL_END<<"\n";
 }
 
 Network::~Network() {

src/NetworkRT.cpp

@@ -1,9 +1,13 @@
 #include <iostream>
 #include <map>
 #include <errno.h>
+#include <string.h> // memcpy
+#include <stdlib.h>
 
+#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 w { dtRT, data_b, l->inputs*l->outputs};
-    Weights b = { dtRT, l->bias_h, 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;
+        Weights power{dtRT, power_b, l->outputs};
-
+        Weights shift{dtRT, mean_b, l->outputs};
-        //make power array of ones
+        Weights scale{dtRT, variance_b, l->outputs};
-        dnnType *power_h = new dnnType[l->outputs];
-        for(int i=0; i<l->outputs; i++) power_h[i] = 1.0f;
-
-        //convert mean
-        for(int i=0; i<l->outputs; i++)
-            l->mean_h[i] = l->mean_h[i] / -sqrt(eps + l->variance_h[i]); 
-        
-        //convert variance
-        for(int i=0; i<l->outputs; 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};
         IScaleLayer *lRT2 = networkRT->addScale(*lRT->getOutput(0), ScaleMode::kCHANNEL, 
                     shift, scale, power);
         checkNULL(lRT2);
 
-        Weights shift2{dtRT, l->bias_h, l->outputs};
+        Weights shift2{dtRT, bias_b, l->outputs};
-        Weights scale2{dtRT, l->scales_h, l->outputs};
+        Weights scale2{dtRT, scales_b, l->outputs};
         IScaleLayer *lRT3 = networkRT->addScale(*lRT2->getOutput(0), ScaleMode::kCHANNEL, 
                     shift2, scale2, power);
         checkNULL(lRT3);

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<size) {    
+        output[i] = __float2half(input[i]);
+    }
+}
+
+
+void float2half(float* srcData, __half *dstData, int size, const cudaStream_t stream)
+{
+    int blocks = (size+255)/256;
+    int threads = 256;
+    
+    float2half_device<<<blocks, threads, 0, stream>>>(srcData, dstData, size);
+    cudaDeviceSynchronize();
+}

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];

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