用openCL/C实现矩阵乘法

Question

我了解矩阵乘法的理论，我对这个特定的内核实现只有两个问题：

作为参考，num_rows = 32。矩阵 B (b_mat) 之前已被另一个内核转置，所以据我了解，我们将行向量点在一起。

1) 为什么我们需要使用参数“vectors_per_row”以及内部循环？我以为我们可以只做 sum += dot(row of A, row of B)，看起来这个参数正在将行分成更小的部分（为什么？）。

2) 我不明白 a_mat 和 b_mat 的地址偏移量，即 a_mat += start; b_mat += 开始*4；

__kernel void matrix_mult(__global float4 *a_mat,
   __global float4 *b_mat, __global float *c_mat) {
   float sum;
   int num_rows = get_global_size(0);
   int vectors_per_row = num_rows/4;
   int start = get_global_id(0) * vectors_per_row;    
   a_mat += start;                                          
   c_mat += start*4;                                  
   for(int i=0; i<num_rows; i++) {             
      sum = 0.0f;                                      
      for(int j=0; j<vectors_per_row; j++) {   
         sum += dot(a_mat[j],                  
                b_mat[i*vectors_per_row + j]); 
      }                                     
      c_mat[i] = sum;                           
   }                                        
}

Answer 1

1. 您的矩阵由一组 float4 组成。 Flaoa4 是 4 个浮点数的向量。这就是 4 的由来。 Dot 仅适用于内置类型，因此您必须在 float4 上执行此操作。

2. c_mat 是 float 类型，这就是为什么它有 start*4 而 a_mat 有 start。偏移量是因为代码被分成几个（可能数百个）线程。每个线程只计算乘法运算的一小部分。 start 只是线程开始计算的地方。这就是 get_global_id(0) 的用途。它本质上是获取你的线程 ID。从技术上讲，它是第一个维度的线程索引，但您似乎只有一个线程维度，所以在这里您可以将其视为线程 id。