0

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

作为参考,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;                           
   }                                        
}
4

1 回答 1

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

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

于 2012-11-04T00:38:17.637 回答