所谓矩阵的乘法,就是现有两个矩阵,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];
}
}
}