卷積轉換爲矩陣運算中填充數的計算-GEMM

時間  2019-11-06
標籤 轉換 矩陣 運算 填充 計算 gemm 欄目 應用數學 简体版
原文   原文鏈接

    背景：最近在寫一個基於opencl的正向神經網絡框架，項目地址 https://github.com/aktiger/YoloOCLInference ，我從這裏https://github.com/pengdada/YoloOCLInference fork了一個基本的腳手架，可是原始的項目只支持windows的版本，首先把它移植到linux下，因爲須要支持resnet18，還缺乏7*7的卷積，須要本身搞一個，在搞以前，先對3*3的卷積計算進行了梳理，後面7*7的也就瓜熟蒂落。基於opencl作的目的也是爲了可以上嵌入式設備，不想一直生活在服務器的世界裏，因此該造的輪子仍是要本身造。 php

    雖然學術圈仍是工業界都在說卷積神經網絡，可是到了底層要麼是轉換爲矩陣的運算，要麼是轉換爲頻域上的計算。要將卷積操做轉換爲矩陣乘積的第一步是要作img2col操做，以下圖，想詳細瞭解看這裏：https://petewarden.com/2015/04/20/why-gemm-is-at-the-heart-of-deep-learning/  linux

image

 

在執行convertImageTocolumn操做的時候，傳入的原始圖片Img的大小爲3*416*416，填充數爲1，stride爲1，若是填充的塊都爲1， 那麼將Img變成column形式後，裏面會有多少個零呢？答案是14964個，這是程序給出的結果，那麼這個是怎麼來的呢？首先將14964作質數分解：14963=43*3*2*2*29 ，這裏咱們發現有3，那麼這個3能夠看作是三個通道，這樣就只須要看14963=3*4988, 一個通道上有4988個零是怎麼來的。以下圖所示，若是通道大小是3*3，卷積核大小是3*3，填充爲1，步長（stride）爲1，那麼填充後的大小爲5*5，如今要用3*3的卷積在其上進行塊轉換，咱們知道轉換成塊後的大小將是原來大小的9倍。咱們來計算5*5的通道上進行塊轉換後，裏面具備0的數量，首先4個角上，每一個角上得到的填充數爲5，共20個，每一條邊上得到的填充數量爲3，共4條邊，共12個填充，總共爲1*3*4+20 = 32個。 以此類推，若是單個通道的大小變爲416*416，卷積核依然爲3*3，那麼每條邊上得到的填充數量爲(416-3+1)*3 *4+ 4*5 = 4988個，若是通道數量爲3，那麼就爲3*4988=14964。 由此能夠獲得以下的計算公式： git

假設，通道數爲C，通道高度爲H，通道寬度爲W，卷積核大小爲K，填充爲1，步長爲1，假設H=W，那麼得到填充數量爲：3*[(W-K+1)*3*4 + 5*4]. github

 

662311451174692636

程序驗證：ubuntu

image

 

image

 

image

 

將卷積核、全部輸入通道出了填充以外的值都所有設置成1，填充的值設置爲0，對3*416*416的輸入，利用GEMM進行卷積計算，卷積核大小爲3*3，步長和windows

填充都爲1，輸出通道數量爲16，計算完後，獲得不一樣值的分佈以下，結果爲12的數量爲64，值爲18的數量爲26496，值爲27的爲2742336。 能夠知道，輸出服務器

元素的個數爲：16*416*416，即有16個通道，每個通道的大小爲416*416。網絡

image

       如今咱們來分析上面每個值數量的來由： 首先分析，結果爲12的數量64，因爲有16個輸出通道，那麼至關於每個輸出通道有4個12的值，這正好對應於在每一個輸入通道的4個角上卷積後的結果，由於四個角上，有5個值是被填充的，還剩下4個數爲1，乘以輸入通道數3，獲得的數值正好爲12。 而後分析值爲18的數量26496，一樣道理，因爲輸出通道有16，至關於每個輸出通道上有26496/16=1656個值爲18的數。對應於每個輸入通道邊上的卷積（不包括4個角）的數量爲（416-3+1）*4=1656，而每個這樣的卷積中，有3個數是被填充的，3個通道就是有9個數是被填充的，原本是27個1，可是這裏有9個1被填充爲0，因此最後的數值就位18。最後分析值爲27的數量：2742336，一樣，因爲有16個輸出通道，那麼每個通道27的數量爲2742336/16=171396個，在輸入通道上除了邊上和和4個角上的卷積（內部卷積，每一個位置都是1）的數量爲414*414，正好等於171396個。 框架

 

能夠用這個在線質數分解器：http://www.atool.org/quality_factor.phpide

計算im2col輸出元素個數的公式爲： in_ch * k_size * k_size * out_w * out * out_h

其中：

in_ch:  輸入通道大小

k_size:  卷積核大小

out_w: 在輸入圖片上作卷積後，輸出通道的寬，計算公式爲：out_w = (in_w + 2*pad  - k_size )/stride + 1

out_h： 在輸入圖片上作卷積後，輸出通道的高， 計算公式爲： out_h = (in_h + 2*pad –k_size )/stride + 1

下面是運行3*3卷積和7*7卷積的日誌：

 

/home/ubuntu/zhangchao/cvs/YoloOCLInference/cmake-build-debug/test/test
Running 2 test cases...
CL_COMPUTE DEVICES: 2
CL_DEVICE_ID: 0x2594750
CL_DEVICE_NAME:: Tesla K40m
CL_DEVICE_VENDOR:: NVIDIA Corporation
CL_DRIVER_VERSION:: 375.26
CL_DEVICE_VERSION:: OpenCL 1.2 CUDA
CL_DEVICE_OPENCL_C_VERSION:: OpenCL C 1.2 
CL_DEVICE_TYPE::CL_DEVICE_TYPE_GPU
CL_DEVICE_MAX_COMPUTE_UNITS: 15
clCreateProgramWithSource success
clGetProgramBuildInfo() success
buildProgram kernels success
Kernel No: 1, name - image2columarray3x3
buildProgram kernels success
Kernel No: 2, name - image2columarray1x1
buildProgram kernels success
Kernel No: 3, name - resetarray
buildProgram kernels success
Kernel No: 4, name - normalizearray
buildProgram kernels success
Kernel No: 5, name - scalebias
buildProgram kernels success
Kernel No: 6, name - addbias
buildProgram kernels success
Kernel No: 7, name - scaleaddbias
buildProgram kernels success
Kernel No: 8, name - normscaleaddbias
buildProgram kernels success
Kernel No: 9, name - leakyactivatearray
buildProgram kernels success
Kernel No: 10, name - linearactivatearray
buildProgram kernels success
Kernel No: 11, name - flattenarray
buildProgram kernels success
Kernel No: 12, name - softmax
buildProgram kernels success
Kernel No: 13, name - maxpool
buildProgram kernels success
Kernel No: 14, name - image2columarray7x7
Number of kernel Arguments : 11 image2columarray3x3 
Number of kernel Arguments : 11 image2columarray1x1 
Number of kernel Arguments : 7 normalizearray 
Number of kernel Arguments : 4 scalebias 
Number of kernel Arguments : 4 addbias 
Number of kernel Arguments : 5 scaleaddbias 
Number of kernel Arguments : 7 normscaleaddbias 
Number of kernel Arguments : 4 leakyactivatearray 
Number of kernel Arguments : 4 linearactivatearray 
Number of kernel Arguments : 6 flattenarray 
Number of kernel Arguments : 7 softmax 
Number of kernel Arguments : 9 maxpool 
Number of kernel Arguments : 3 resetarray 
Number of kernel Arguments : 11 image2columarray7x7 
In bufImg_before7x7.bin. Total count is : 519168;   Zero count is :0. Percent is: 0
In bufImg9x_before7x7.bin. Total count is : 24952368;   Zero count is :0. Percent is: 0
In buf_out_before7x7.bin. Total count is : 2715904;   Zero count is :2715904. Percent is: 1
In weights_gpu7x7.bin. Total count is : 2352;   Zero count is :0. Percent is: 0
Total kernel time was {-42898.564} msecs - image2columarray7x7 
CL_Status is not CL_SUCCESS
get_local_size(0):8, get_num_groups(0): 63655, get_global_id(0): 509239
In bufImg9x_after.bin. Total count is : 24952368;   Zero count is :34596. Percent is: 0.00138648
Total kernel time was { 0.00} msecs - ComputeGEMM() 
In buf_out_after7x7.bin. Total count is : 2715904;   Zero count is :0. Percent is: 0
data_img zero count is: 0. data_img_count: 519168
data_in zero count is: 0
data_out zero count is: 2768896
kernel_weights zero count is: 0
CL_COMPUTE DEVICES: 2
CL_DEVICE_ID: 0x2594750
CL_DEVICE_NAME:: Tesla K40m
CL_DEVICE_VENDOR:: NVIDIA Corporation
CL_DRIVER_VERSION:: 375.26
CL_DEVICE_VERSION:: OpenCL 1.2 CUDA
CL_DEVICE_OPENCL_C_VERSION:: OpenCL C 1.2 
CL_DEVICE_TYPE::CL_DEVICE_TYPE_GPU
CL_DEVICE_MAX_COMPUTE_UNITS: 15
clCreateProgramWithSource success
clGetProgramBuildInfo() success
buildProgram kernels success
Kernel No: 1, name - image2columarray3x3
buildProgram kernels success
Kernel No: 2, name - image2columarray1x1
buildProgram kernels success
Kernel No: 3, name - resetarray
buildProgram kernels success
Kernel No: 4, name - normalizearray
buildProgram kernels success
Kernel No: 5, name - scalebias
buildProgram kernels success
Kernel No: 6, name - addbias
buildProgram kernels success
Kernel No: 7, name - scaleaddbias
buildProgram kernels success
Kernel No: 8, name - normscaleaddbias
buildProgram kernels success
Kernel No: 9, name - leakyactivatearray
buildProgram kernels success
Kernel No: 10, name - linearactivatearray
buildProgram kernels success
Kernel No: 11, name - flattenarray
buildProgram kernels success
Kernel No: 12, name - softmax
buildProgram kernels success
Kernel No: 13, name - maxpool
buildProgram kernels success
Kernel No: 14, name - image2columarray7x7
Number of kernel Arguments : 11 image2columarray3x3 
Number of kernel Arguments : 11 image2columarray1x1 
Number of kernel Arguments : 7 normalizearray 
Number of kernel Arguments : 4 scalebias 
Number of kernel Arguments : 4 addbias 
Number of kernel Arguments : 5 scaleaddbias 
Number of kernel Arguments : 7 normscaleaddbias 
Number of kernel Arguments : 4 leakyactivatearray 
Number of kernel Arguments : 4 linearactivatearray 
Number of kernel Arguments : 6 flattenarray 
Number of kernel Arguments : 7 softmax 
Number of kernel Arguments : 9 maxpool 
Number of kernel Arguments : 3 resetarray 
Number of kernel Arguments : 11 image2columarray7x7 
In bufImg_before.bin. Total count is : 519168;   Zero count is :0. Percent is: 0
In bufImg9x_before.bin. Total count is : 4672512;   Zero count is :0. Percent is: 0
In databuf_out_before.bin. Total count is : 2768896;   Zero count is :2768896. Percent is: 1
In weights_gpu.bin. Total count is : 432;   Zero count is :0. Percent is: 0
Total kernel time was {-42882.983} msecs - image2columarray3x3 
CL_Status is not CL_SUCCESS
get_local_size(0):8, get_num_groups(0): 64897, get_global_id(0): 519175
In bufImg9x_after.bin. Total count is : 4672512;   Zero count is :14964. Percent is: 0.00320256
Total kernel time was { 0.00} msecs - ComputeGEMM() 
In databuf_out_after.bin. Total count is : 2768896;   Zero count is :0. Percent is: 0

*** No errors detected

Process finished with exit code 0
相關文章
相關標籤/搜索
每日一句
    每一个你不满意的现在，都有一个你没有努力的曾经。
本站公眾號
   歡迎關注本站公眾號,獲取更多信息
聯繫我們 最近搜索 最新文章 沪ICP备13005482号-10 MyBatis教程 SQL 教程 MySQL教程 Java 教程 Thymeleaf 教程 Hibernate教程 Spring教程 Redis教程