优化矩阵乘法 OpenCL - 目的：学习如何管理内存

Question

我是 OpenCL 的新手，并试图了解如何优化矩阵乘法以熟悉各种范式。这是当前代码。 如果我将矩阵 A 和 B 相乘。我在私有内存中分配 A 行作为开始（因为每个工作项都使用它），并在本地内存中分配 B 列（因为每个工作组都使用它）。

1) 代码当前不正确，不幸的是我正在努力研究如何使用本地工作 ID 来获取正确的代码，但我找不到我的错误？我以http://www.cs.bris.ac.uk/home/simonm/workshops/OpenCL_lecture3.pdf 为基础，但是（幻灯片 27）这似乎是错误的，因为他们没有在内部循环中使用 loc_size）

2) 您对此代码还有什么其他优化建议吗？

 __kernel void mmul(
  __global int* C,
 __global int* A, 
 __global int* B,
   const int rA, 
   const int rB,
   const int cC, 
   __local char* local_mem) 
{ 
   int k,ty; 
   int tx = get_global_id(0); 
   int loctx = get_local_id(0); 
   int loc_size = get_local_size(0);
   int value = 0 ;
   int tmp_array[1000]; 
   for(k=0;k<rB;k++) { 
     tmp_array[k] = A[tx * cA + k] ;
   } 
   for (ty=0 ; ty < cC ; ty++) { \n" \
     for (k = loctx ; k < rB ; k+=loc_size) { 
         local_mem[k] = B[ty + k * cC] ;
     }
      barrier(CLK_LOCAL_MEM_FENCE); 
       value = 0 ; 
       for(k=0;k<rB;k+=1) {
           int i = loctx + k*loc_size;
           value += tmp_array[k] * local_mem[i]; 
     } 
   C[ty + (tx * cC)] = value; 
 } 
} 

我在哪里设置全局和本地工作项如下

const size_t globalWorkItems[1] = {result_row};
const size_t localWorkItems[1] = {(size_t)local_wi_size};

local_wi_size 是 result_row/计算单元数（这样 result_row % 计算单元 == 0）

Answer 1

您的代码非常接近，但是对本地内存数组的索引实际上比您想象的要简单。您在私有内存中有一行，在本地内存中有一个列，您需要计算这两个向量的点积。你只需要将row[k]*col[k] 相加，对于k = 0 最高可达N-1：

for(k=0;k<rB;k+=1) {
    value += tmp_array[k] * local_mem[k]; 
}

实际上，在您使用的幻灯片中给出的示例解决方案中也存在第二个更微妙的错误。由于您在循环中读取和写入本地内存，因此您实际上需要 两个 屏障，以确保在迭代 i 时写入本地内存的工作项不会覆盖正在被其他执行迭代的工作项读取i-1。

因此，您的内核的完整代码（经过测试和工作）应该如下所示：

__kernel void mmul(
  __global int* C,
  __global int* A,
  __global int* B,
     const int rA,
     const int rB,
     const int cC,
  __local char* local_mem)
{
  int k,ty;
  int tx = get_global_id(0);
  int loctx = get_local_id(0);
  int loc_size = get_local_size(0);
  int value = 0;
  int tmp_array[1000];
  for(k=0;k<rB;k++) {
    tmp_array[k] = A[tx * cA + k] ;
  }
  for (ty=0 ; ty < cC ; ty++) {

    for (k = loctx ; k < rB ; k+=loc_size) {
      local_mem[k] = B[ty + k * cC];
    }
    barrier(CLK_LOCAL_MEM_FENCE); // First barrier to ensure writes have finished

    value = 0;
    for(k=0;k<rB;k+=1) {
      value += tmp_array[k] * local_mem[k];
    }
    C[ty + (tx * cC)] = value;

    barrier(CLK_LOCAL_MEM_FENCE); // Second barrier to ensure reads have finished
  }
}

您可以在HandsOnOpenCL GitHub page 上找到与您正在查看的幻灯片配套的全套练习和解决方案。还有一组来自同一教程的更完整的幻灯片here，它继续展示了一个更加优化的矩阵乘法示例，该示例使用阻塞方法来更好地利用时间和空间局部性。前面提到的缺少屏障错误已在示例解决方案代码中修复，但尚未在幻灯片中修复。