0

行列の乗算の理論を理解しています。この特定のカーネルの実装について2つの質問があります。

参考までに、num_rows =32です。行列B（b_mat）は、以前に別のカーネルによって置き換えられたため、私が理解しているように、行ベクトルをドットでつなぎ合わせています。

1）パラメータ "vectors_per_row"を使用する必要があるのはなぜですか？したがって、内部ループを使用する必要がありますか？sum + = dot（Aの行、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を取得します。技術的には最初のディメンションのスレッドインデックスですが、スレッドディメンションは1つしかないように見えるため、ここではスレッドIDと考えることができます。

于 2012-11-04T00:38:17.637 に答える