卷积转换为矩阵运算中填充数的计算-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教程