2

CUDA で行列の乗算に使用するコードでは、正方行列と非正方行列の両方を乗算できますが、幅と高さの両方がブロックサイズの倍数でなければなりません。

たとえば、[3][6] * [6][3] (blocksize=3 を使用) を乗算することはできますが、[3][2]*[2][3] を乗算することはできません。

誰もそれを行う方法を知っていますか? これは私のカーネルです:

#include <stdio.h>

#include <limits.h>

#include <stdlib.h>
#define blocksize 3
#define HM (1*blocksize) 
#define WM (2*blocksize) 
#define WN (1*blocksize)
#define HN WM 
#define WP WN   
#define HP HM  
#define PTH WM
#define PTW HM

__global__ void nonsquare(float*M, float*N, float*P, int uWM,int uWN)

{
__shared__ float MS[blocksize][blocksize];
__shared__ float NS[blocksize][blocksize];


int tx=threadIdx.x, ty=threadIdx.y, bx=blockIdx.x, by=blockIdx.y;
int rowM=ty+by*blocksize;
int colN=tx+bx*blocksize;
float Pvalue=0;


for(int m=0; m< uWM/blocksize;++m){
    MS[ty][tx]=M[rowM*uWM+(m*blocksize+tx)];
    NS[ty][tx]=M[colN + uWN*(m*blocksize+ty)];
    __syncthreads();

    for(int k=0;k<blocksize;k++)
        Pvalue+=MS[ty][k]*NS[k][tx];
    __syncthreads();
    P[rowM*WP+colN]=Pvalue;
     }
    }

前もって感謝します！

4

2 に答える 2

5

最も簡単な方法は、最後のブロックをゼロで埋めることだと思います。

for(int m=0; m< uWM/blocksize;++m){
    colM = m*blocksize+tx;
    rowN = m*blocksize+ty;
    if (rowM > uWN || rowN > uWM || colM > uWM || colN > uWN) {
        MS[ty][tx]=0.;
        NS[ty][tx]=0.;
    } else {
        MS[ty][tx]=M[rowM*uWM+colM];
        NS[ty][tx]=N[colN + uWN*rowN];
    }

プラスかマイナス。（そのNS行はMではなくNを参照する必要がありますよね？）

しかし、可能であれば既存の調整済みライブラリを使用することを提唱しているのは私だけのようです。自分で作成する代わりに、CUBLASまたはMAGMAを使用してみませんか？それらは高速で、何百人ものユーザーによってテストされています。

于 2011-03-28T15:28:07.467 に答える
2

ここでの基本的なパフォーマンス要件は、共有メモリ「タイル」の 1 番目または 2 番目の次元のいずれかが 16 の倍数であるということです。これは、最適なグローバル メモリ帯域幅 (つまり、ハーフ ワープ合体トランザクション) を達成するために歴史的に必要なものです。タイルの 1 番目または 2 番目の次元のどちらにするかは、行列が列優先順または行優先順のどちらで格納されているかによって決まります。共有メモリ タイルが正方形である必要があることは言うまでもありませんが、ストレージの最初の次元 (BLAS 表記の LDA) が 16 の倍数に丸められることだけが必要です。

タイルの次元をテンプレート引数として使用してカーネルを簡単にテンプレート化し、行列の次元に応じていくつかのバージョンをインスタンス化できます。特定のアーキテクチャに対して、占有率と命令レベルの並列処理のバランスをとる最適なタイル ディメンションが存在する可能性があります。これを解決する「賢い」方法は、おそらく行列の乗算を 2 つの操作に分解することです。1 つ目は最適なタイル サイズで大部分の作業を行い、2 つ目は残りの列の異なるサイズで行います。製品が完成した後、結果がホスト メモリに直接戻る場合、GPU カーネルとオーバーラップする最適化された BLAS を使用して、ホスト上で 2 番目の操作を実行するのが最適な場合があります。これは、UTK Magma ライブラリのルーチンの多くが使用するアプローチです。

于 2011-03-29T10:29:21.673 に答える