2. 程式人生 > >讓Faster R-CNN支援TX1的fp16(half float, float16)特性

讓Faster R-CNN支援TX1的fp16(half float, float16)特性

阿新 • • 發佈：2019-01-15

為什麼要用float16？

一句話，float16的運算速度是float32運算速度的2倍。nVidia說的，資料不對不要怪我，(≖ ‿ ≖)✧

在說一下缺點，也是一句話，精度損失了。（臥槽，廢話，(*´Д｀*)）

會損失多少呢？

IEEE754（wiki）這裡描述各種float的規則，這裡鄙視一下百度百科。 눈_눈

float32: 負 -3.4028235E+38 到 -1.401298E-45，正 1.401298E-45 到 3.4028235E+38

float16: 半精度佔2個位元組，1位符號位，5位階碼，10位尾陣列成，精度為0.001，所以它的範圍，自己算 ，(´・ω・`) ，不會就去問老師(｡◕ˇ∀ˇ◕）。

看來還是可以和float32比一比的嘛，（喂喂，這麼果斷得出結論真的好嗎？）

既然那麼厲害， 反正拿到了TX1，那就試試呀，（臥槽，真的不打算解釋了麼）！搞一個超大的矩陣乘，然而事實是cuda目前也只提供float16的矩陣乘運算，_(:3 」∠)_

先貼個程式碼，請直接略過，程式碼有點長，有點亂，請隨便介意，因為不是我寫的，(｡◕ˇ∀ˇ◕）

  1. #include <algorithm>
  2. #include <iostream>
  3. #include <time.h>
  4. #include <cublas.h>
  5. #include <cublas_v2.h>
  6. #include <assert.h>
  7. #include <stdio.h>
  8. #include <cuda_fp16.h>
  9. usingnamespace std;  
  10. #define  IDX2C(i,j,leading) (((j)*(leading))+(i))
  11. typedefstruct _data *PDATA;  
  12. typedefstruct _data  
  13. {  
  14.     int _rows;  
  15.     int _cols;  
  16.     float *data;  
  17. } Data;  
  18. typedefstruct _hdata *PHDATA;  
  19. typedefstruct
     _hdata  
  20. {  
  21.     int _rows;  
  22.     int _cols;  
  23.     half *data;  
  24. } HData;  
  25. void free_mat(PDATA mat)  
  26. {  
  27.     free(mat->data);  
  28.     free(mat);  
  29. }  
  30. typedef unsigned short uint16_t;  
  31. typedef unsigned int uint32_t;  
  32. half uint16_as_fp16 (uint16_t a)  
  33. {  
  34.     half res;  
  35. #if defined (__cplusplus)
  36.     memcpy (&res, &a, sizeof (res));  
  37. #else /* __cplusplus */
  38.     volatileunion {  
  39.         half f;  
  40.         uint16_t i;  
  41.     } cvt;  
  42.     cvt.i = a;  
  43.     res = cvt.f;  
  44. #endif /* __cplusplus */
  45.     return res;  
  46. }  
  47. uint32_t fp32_as_uint32 (float a)  
  48. {  
  49.     uint32_t res;  
  50. #if defined (__cplusplus)
  51.     memcpy (&res, &a, sizeof (res));  
  52. #else /* __cplusplus */
  53.     volatileunion {  
  54.         float f;  
  55.         uint32_t i;  
  56.     } cvt;  
  57.     cvt.f = a;  
  58.     res = cvt.i;  
  59. #endif /* __cplusplus */
  60.     return res;  
  61. }  
  62. /* host version of device function __float2half_rn() */
  63. half float2half_rn (float a)  
  64. {  
  65.     uint32_t ia = fp32_as_uint32 (a);  
  66.     uint16_t ir;  
  67.     ir = (ia >> 16) & 0x8000;  
  68.     if ((ia & 0x7f800000) == 0x7f800000) {  
  69.         if ((ia & 0x7fffffff) == 0x7f800000) {  
  70.             ir |= 0x7c00; /* infinity */
  71.         } else {  
  72.             ir = 0x7fff; /* canonical NaN */
  73.         }  
  74.     } elseif ((ia & 0x7f800000) >= 0x33000000) {  
  75.         int shift = (int)((ia >> 23) & 0xff) - 127;  
  76.         if (shift > 15) {  
  77.             ir |= 0x7c00; /* infinity */
  78.         } else {  
  79.             ia = (ia & 0x007fffff) | 0x00800000; /* extract mantissa */
  80.             if (shift < -14) { /* denormal */
  81.                 ir |= ia >> (-1 - shift);  
  82.                 ia = ia << (32 - (-1 - shift));  
  83.             } else { /* normal */
  84.                 ir |= ia >> (24 - 11);  
  85.                 ia = ia << (32 - (24 - 11));  
  86.                 ir = ir + ((14 + shift) << 10);  
  87.             }  
  88.             /* IEEE-754 round to nearest of even */
  89.             if ((ia > 0x80000000) || ((ia == 0x80000000) && (ir & 1))) {  
  90.                 ir++;  
  91.             }  
  92.         }  
  93.     }  
  94.     return uint16_as_fp16 (ir);  
  95. }  
  96. PHDATA mat_product(PHDATA mat1,PHDATA mat2)  
  97. {  
  98.     if(mat1->_cols!=mat2->_rows)  
  99.     {  
  100.         printf("this is not right\n");  
  101.             return NULL;  
  102.     }  
  103.     PHDATA mat3=new HData;  
  104.     mat3->data=(half *)malloc(sizeof(half)*(mat1->_rows)*(mat2->_cols));  
  105.     mat3->_rows=mat1->_rows;  
  106.     mat3->_cols=mat2->_cols;  
  107.     /* 
  108.      *INIT the matrix we want calculate  
  109.      * col primary 
  110.      */
  111.     {  
  112.         half *d_a,*d_b,*d_c;  
  113.         cublasInit();  
  114.         cublasAlloc((mat1->_cols)*(mat1->_rows),sizeof(half),(void **)&d_a);  
  115.         cublasAlloc((mat2->_cols)*(mat2->_rows),sizeof(half),(void **)&d_b);  
  116.         cublasAlloc((mat3->_rows)*(mat3->_cols),sizeof(half),(void **)&d_c);  
  117.         cudaMemcpy(d_a,mat1->data,sizeof(half)*(mat1->_cols)*(mat1->_rows),cudaMemcpyHostToDevice);  
  118.         cudaMemcpy(d_b,mat2->data,sizeof(half)*(mat2->_rows)*(mat2->_cols),cudaMemcpyHostToDevice);  
  119.         cublasHandle_t handle;  
  120.         cublasCreate(&handle);  
  121.         half alpha=float2half_rn(float(1.0));  
  122.         half beta=float2half_rn(float(0.0));  
  123.         cudaEvent_t start,stop;  
  124.  <span style="white-space:pre">   </span>cudaEventCreate(&start);  
  125.  <span style="white-space:pre">   </span>cudaEventCreate(&stop);  
  126.  <span style="white-space:pre">   </span>cudaEventRecord(start,0);  
  127.         cublasHgemm(handle,CUBLAS_OP_N,CUBLAS_OP_N,mat1->_rows,mat2->_cols,  
  128.             mat2->_rows,&alpha,d_a,mat1->_rows,d_b,mat2->_rows,&beta,d_c,mat1->_rows);  
  129.  <span style="white-space:pre">   </span>cudaEventRecord(stop,0);  
  130.  <span style="white-space:pre">   </span>cudaEventSynchronize(stop);  
  131.  <span style="white-space:pre">   </span>float ela=0;  
  132.  <span style="white-space:pre">   </span>cudaEventElapsedTime(&ela,start,stop);  
  133.  <span style="white-space:pre">   </span>cout<<"GPU: "<<ela<<"ms"<<endl;  
  134.         cudaMemcpy(mat3->data,d_c,sizeof(half)*(mat3->_rows)*(mat3->_cols),cudaMemcpyDeviceToHost);  
  135.         cublasFree(d_a);  
  136.         cublasFree(d_b);  
  137.         cublasFree(d_c);  
  138.         cublasShutdown();  
  139.     }  
  140.     /* need to trans the mat3*/
  141.     return mat3;  
  142. }  
  143. void ele_mat_show(PDATA mat)  
  144. {  
  145.     for (int i=0;i<mat->_rows;i++){  
  146.         for (int j=0;j<mat->_cols;j++){  
  147.             cout<<mat->data[IDX2C(i,j,mat->_rows)]<<"\t";  
  148.         }  
  149.         cout<<endl;  
  150.     }  
  151. }  
  152. float myrand()  
  153. {  
  154.     return rand()%10;  
  155. }  
  156. int main()  
  157. {  
  158.     //clock_t start,end;
  159. #if 0
  160.     for (int i=0;i<M*N;i++)  
  161.     {  
  162.         cout<<c[i]<<"\t";  
  163.     }  
  164.     cout<<endl;  
  165. #endif
  166.     PDATA mat1,mat2,mat3;  
  167.     PHDATA mat4,mat5,mat6;  
  168.     /* remember to initialize the point*/
  169.     mat1=(PDATA)malloc(sizeof(Data));  
  170.     mat2=(PDATA)malloc(sizeof(Data));  
  171.     mat3=(PDATA)malloc(sizeof(Data));  
  172.     mat4=(PHDATA)malloc(sizeof(HData));  
  173.     mat5=(PHDATA)malloc(sizeof(HData));  
  174.     mat6=(PHDATA)malloc(sizeof(HData));  
  175.     mat1->_rows=5000;  
  176.     mat1->_cols=50000;  
  177.     mat4->_rows=5000;  
  178.     mat4->_cols=50000;  
  179.     mat1->data=(float *)malloc(sizeof(float)*mat1->_rows*mat1->_cols);  
  180.     mat4->data=(half *)malloc(sizeof(half)*mat1->_rows*mat1->_cols);  
  181.     for (int i=0;i<mat1->_rows;i++)  
  182.         for (int j=0;j<mat1->_cols;j++)  
  183.             mat1->data[IDX2C(i,j,mat1->_rows)]=i+j;  
  184.     for (int i=0;i<mat1->_rows;i++)  
  185.         for (int j=0;j<mat1->_cols;j++)  
  186.             mat4->data[IDX2C(i,j,mat1->_rows)]=float2half_rn(mat1->data[IDX2C(i,j,mat1->_rows)]);  
  187.     mat2->_rows=50000;  
  188.     mat2->_cols=2000;  
  189.     mat5->_rows=50000;  
  190.     mat5->_cols=2000;  
  191.     mat2->data=(float *)malloc(sizeof(float)*mat2->_rows*mat2->_cols);  
  192.     mat5->data=(half *)malloc(sizeof(half)*mat2->_rows*mat2->_cols);  
  193.     for (int i=0;i<mat2->_rows;i++)  
  194.         for (int j=0;j<mat2->_cols;j++)  
  195.             mat2->data[IDX2C(i,j,mat2->_rows)]=i+j;  
  196.     for (int i=0;i<mat2->_rows;i++)  
  197.         for (int j=0;j<mat2->_cols;j++)  
  198.             mat5->data[IDX2C(i,j,mat2->_rows)]=float2half_rn(mat2->data[IDX2C(i,j,mat2->_rows)]);  
  199.     mat6=mat_product(mat4,mat5);  
  200.     return 0;  
  201. }  
漫長的等待之後，跑了完了， 結果是什麼呢？ (´・ω・`) 

我擦咧， (╯‵□′)╯︵┻━┻還真尼瑪快了兩倍，nVidia你贏了。

那如此一來，就得開始幹活啦，讓faster rcnn支援TX1的半精度運算。

怎麼做

1.環境：cuda7.5，python，TX1等等 2.需要caffe支援半精運算 3.新增faster rcnn新增的層：roi_pooling_layer和smooth_L1_loss_layer 4.測試（做好後將會在我的部落格中公佈結果）

目前情況

1.Faster R-CNN在TX1上執行。 2.讓mnist在TX1上跑半精運算。 3.正在重寫roi_pooling_layer和smooth_L1_loss_layer

相關推薦

讓Faster R-CNN支援TX1的fp16(half float, float16)特性

為什麼要用float16？ 一句話，float16的運算速度是float32運算速度的2倍。nVidia說的，資料不對不要怪我，(≖ ‿ ≖)✧ 在說一下缺點，也是一句話，精度損失了。（臥槽，廢話，(*´Д｀*)） 會損失多少呢？ IEEE754（wiki）這裡描述各種f

tensorflow object detection faster r-cnn 中keep_aspect_ratio_resizer是什麽意思

ng- 最小 圖片 sta fault overflow cti hub .com 如果小夥伴的英語能力強可以直接閱讀這裏：https://stackoverflow.com/questions/45137835/what-the-impact-of-differe

Faster R-CNN技巧

工作 並行 labels ast 個數 draw 默認 物體 ray 1.盡量控制圖片大小在1024以內，不然顯存會爆炸。 2.盡量使用多GPU並行工作，訓練下降速度快。 3.當需要被檢測的單張圖片裏物體太多時，記得修改Region_proposals的個數 4.測試的時候

Faster R-CNN：詳解目標檢測的實現過程

最大的 中心 width 小數據 等等 eat tar 優先 博文 本文詳細解釋了 Faster R-CNN 的網絡架構和工作流，一步步帶領讀者理解目標檢測的工作原理，作者本人也提供了 Luminoth 實現，供大家參考。 Luminoth 實現：h

純C++版500VIP源碼下載的Faster R-CNN（通過caffe自定義RPN層實現）

方便 預測 大致 ole test cto oop 可執行文件 names 這裏500VIP源碼下載 dsluntan.com 介紹的是通過添加自定義層（RPN層）代替python層，實現c++版的Faster R-CNN，因為去掉python了，所以部署時不會因為牽扯到p

論文閱讀筆記（六）Faster R-CNN: Towards Real-Time Object Detection with Region Proposal Networks

采樣 分享 最終 產生 pre 運算 減少 att 我們 作者：Shaoqing Ren, Kaiming He, Ross Girshick, and Jian SunSPPnet、Fast R-CNN等目標檢測算法已經大幅降低了目標檢測網絡的運行時間。可是盡管如此，仍然

【Faster RCNN】《Faster R-CNN: Towards Real-Time Object Detection with Region Proposal Networks》

NIPS-2015 NIPS，全稱神經資訊處理系統大會(Conference and Workshop on Neural Information Processing Systems)，是一個關於機器學習和計算神經科學的國際會議。該會議固定在每年的12月舉行

R-CNN、Fast R-CNN、Faster R-CNN

在介紹Faster R-CNN之前，先來介紹一些前驗知識，為Faster R-CNN做鋪墊。 一、基於Region Proposal（候選區域）的深度學習目標檢測演算法 Region Proposal（候選區域），就是預先找出圖中目標可能出現的位置，通

Faster R-cnn中的RPN網路詳細解釋

作者RPN網路前面的g層借用的是ZF網路，網路相對較淺，不過並不影響後期介紹。 1、首先，輸入圖片大小是 224*224*3（這個3是三個通道，也就是RGB三種） 2、然後第一層的卷積核維度是 7*7*3*96 （所以大家要認識到卷積核都是4維的，在caffe的矩陣計算中都是這麼實現的）； 3、

Faster R-CNN 英文論文翻譯筆記

Faster R-CNN: Towards Real-Time Object Detection with Region Proposal Networks Shaoqing Ren, Kaiming He, Ross Girshick, Jian Sun reference link：ht

Faster R-CNN改進篇（二）： RFCN ● RON

版權宣告：本文為博主原創文章，未經博主允許不得轉載。 https://blog.csdn.net/linolzhang/article/details/75137050 @改進1：RFCN        論文：R-FCN: Object Detecti

Faster R-CNN改進篇（一）： ION ● HyperNet ● MS CNN

版權宣告：本文為博主原創文章，未經博主允許不得轉載。 https://blog.csdn.net/linolzhang/article/details/74159463 一. 源起於Faster        深度學習於目標檢測的里程碑成果，來自於這篇

Faster R-CNN論文及原始碼解讀

R-CNN是目標檢測領域中十分經典的方法，相比於傳統的手工特徵，R-CNN將卷積神經網路引入，用於提取深度特徵，後接一個分類器判決搜尋區域是否包含目標及其置信度，取得了較為準確的檢測結果。Fast R-CNN和Faster R-CNN是R-CNN的升級版本，在準確率和實時性方面都得到了較大提升。在F

Faster R-CNN 論文學習

Faster R-CNN: Towards Real-Time Object Detection with Region Proposal Networks 演算法簡介 Abstract 1. Introduction 2. Relat

一文讀懂目標檢測 R-CNN Fast R-CNN Faster R-CNN YOLO SSD

分享一下我老師大神的人工智慧教程！零基礎，通俗易懂！http://blog.csdn.net/jiangjunshow 也歡迎大家轉載本篇文章。分享知識，造福人民，實現我們中華民族偉大復興！        

執行Keras版本的Faster R-CNN

Keras版本的Faster R-CNN原始碼下載地址：https://github.com/yhenon/keras-frcnn下載以後，用PyCharm開啟（前提是已經安裝了Tensorflow-gpu和Keras），開啟以後可以看到專案的結構： 修改requirements.txt，設定Keras

Faster R-CNN 目標檢測演算法詳細總結分析（two-stage)(深度學習)(NIPS 2015)

論文名稱：《 Faster R-CNN: Towards Real-Time Object Detection with Region Proposal Networks 》 論文下載：https://papers.nips.cc/paper/5638-faster-r-cnn-t

Faster R-CNN 學習筆記

Faster RCNN 學習 Faster R-CNN這篇論文著重解決了這個系統中的三個問題： 1. 如何設計區域生成網路；RPN 2. 如何訓練區域生成網路； 3. 如何讓區域生成網路和Fast RCNN網路共享特徵提取網路。 圖1：Faster RCNN 結構 Fa

譯文：Faster R-CNN

【翻譯】Faster R-CNN: Towards Real-Time Object Detection with Region Proposal Networks Shaoqing Ren, Kaiming He, Ross Girshick, and Jian Sun 摘要:

Faster R-CNN：利用區域提案網路實現實時目標檢測 論文翻譯

Faster R-CNN論文地址:Faster R-CNN Faster R-CNN專案地址:https://github.com/ShaoqingRen/faster_rcnn 摘要 目前最先進的目標檢測網路需要先用區域提案演算法推測目標位置，像SPPnet1和Fast R-CNN2