矩阵乘优化的思考

所谓矩阵的乘法，就是现有两个矩阵，arr和brr。进行相乘运算，即：arr的第一行乘brr的第一列对应位置之后作和，作为结果的第一行第一列。arr的第一行乘brr的第二列对应位置之后作和，作为结果的第一行第二列。。。。一次类推~

没有做任何优化的代码：
__kernel void MatrixMul(__global const float *arr, __global const float *brr, __global float *crr, const int n)
{
    int i = get_global_id(0);
    int j = get_global_id(1);   
    for(int k = 0; k < n; ++k)
    {
        crr[i*n+j] += arr[i*n+k] * brr[k*n+j];
    }
}

增加每个内核的工作量：
对于1000阶的矩阵，如果每个矩阵元素对应一个工作项，就要求有100万个工作项。所以我们下一步的优化是每个工作项计算矩阵的一行，我们每次都是取arr的一行，然后取brr的各个列之后逐个进行计算，之后放在crr里面的。NDRange从二维的变为了一维。如果有四行，也就是四个计算单元，也就是四个工作组，那么工作组的大小就是250.

__kernel void MatrixMul(__global const float *arr, __global const float *brr, __global float *crr, const int n, const int m)
{
    int i = get_global_id(0);
    for(int j = 0; j < m; ++j)
   {
        for(int k = 0; k < n; ++k)
        {
               crr[i*n+j] += arr[i*n+k] * brr[k*n+j];
        }
    }
}
从二维变成了一维的。

把矩阵A中相关的行从全局内存复制到私有内存：
当做矩阵乘法的时候，arr中的一行要跟所有的列相乘之后相加放到对应行对应列上。那么总是有那么一行正在做运算，反复被用。。这个时候就可以把需要的那一行复制到私有内存中。这样在每一次做运算的时候就不用每一次都去全局内存那里面去找。

__kernel void MatrixMul(__global const float *arr, __global const float *brr, __global float *crr, const int n， const int m, const int p)
{
    int i = get_global_id(0);
    float ans[1000+2];
    for(int k = 0; k < p; ++p)
  ｛
　　　　　　ans[k] = arr[i*n+k];
　｝
    for(int j = 0; j < m; ++j)
    {
            for(int k = 0; k < n; ++k)
            {
                   crr[i*n+j] += ans[k] * brr[k*n+j];
            }
    }
}
 

局部内存：
对于arr我们进行了控制，给一个私有内存区存放，再来看一下brr，arr的每一个行要和brr的各个列进行乘加。arr每换新的一行那么brr就会从开始的第一列开始从新获取一遍知道结束。那么，如果是这样的话，我们就可以定义局部内存保留brr的列。
__kernel void MatrixMul(__global const float *arr, __global const float *brr, __global float *crr, const int n,  const int m, const int p, __local float *cnt)
{
    int i = get_global_id(0);
    int iloc = get_local_id(0);
    int nloc = get_local_size(0);   
    float ans[1000];
    for(int k = 0; k < p; ++p)
  ｛
　　　　　　ans[k] = arr[i*n+k];
　｝
   for(int j = 0; j < m; ++j)
    {
            for(int a = iloc; a < p; a = a + nloc)
                     cnt[a] = cnt[a*p+j];
            for(int k = 0; k < n; ++k)
            {
                   crr[i*n+j] += ant[k] * cnt[k];
            }
    }
}

 

 综上所述是所有的优化方法，我们主要考虑的是尽可能多的不让每一次计算都到全局内存中去读取数据，所以我们在不断的发现不变量和一定时期不变量之后把那些不变的部分转化成私有内存或者是局部内存。