背景：最近在寫一個基於opencl的正向神經網絡框架，項目地址 https://github.com/aktiger/YoloOCLInference ，我從這里https://github.com/pengdada/YoloOCLInference fork了一個基本的腳手架，但是原始的項目只支持windows的版本，首先把它移植到linux下，由於需要支持resnet18，還缺少7*7的卷積，需要自己搞一個，在搞之前，先對3*3的卷積計算進行了梳理，后面7*7的也就順理成章。基於opencl做的目的也是為了能夠上嵌入式設備，不想一直生活在服務器的世界里，所以該造的輪子還是要自己造。
雖然學術圈還是工業界都在說卷積神經網絡，但是到了底層要么是轉換為矩陣的運算，要么是轉換為頻域上的計算。要將卷積操作轉換為矩陣乘積的第一步是要做img2col操作，如下圖，想詳細了解看這里：https://petewarden.com/2015/04/20/why-gemm-is-at-the-heart-of-deep-learning/
在執行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。 由此可以得到如下的計算公式：
假設，通道數為C，通道高度為H，通道寬度為W，卷積核大小為K，填充為1，步長為1，假設H=W，那么獲得填充數量為：3*[(W-K+1)*3*4 + 5*4].
程序驗證：
將卷積核、所有輸入通道出了填充之外的值都全部設置成1，填充的值設置為0，對3*416*416的輸入，利用GEMM進行卷積計算，卷積核大小為3*3，步長和
填充都為1，輸出通道數量為16，計算完后，得到不同值的分布如下，結果為12的數量為64，值為18的數量為26496，值為27的為2742336。 可以知道，輸出
元素的個數為：16*416*416，即有16個通道，每一個通道的大小為416*416。
現在我們來分析上面每一個值數量的來由： 首先分析，結果為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.php
計算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