卷积转换为矩阵运算中填充数的计算-GEMM
Posted justinzhang
tags:
篇首语：本文由小常识网(cha138.com)小编为大家整理，主要介绍了卷积转换为矩阵运算中填充数的计算-GEMM相关的知识，希望对你有一定的参考价值。
背景：最近在写一个基于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
以上是关于卷积转换为矩阵运算中填充数的计算-GEMM的主要内容，如果未能解决你的问题，请参考以下文章