4

GPU に最適化された OpenCL カーネルを作成する方法を学習しようとしています。ローカル メモリ内の正方形タイルを使用した行列乗算の例を取り上げました。ただし、numpy.dot() (5 Gflops 、BLAS を使用) と比較して、せいぜい10 倍のスピードアップ ( ~ 50 Gflops ) しか得られませんでした。

速度が 200 倍以上 (>1000 Gflops) 向上した研究を見つけました。 ftp://ftp.u-aizu.ac.jp/u-aizu/doc/Tech-Report/2012/2012-002.pdf 何が間違っているのか、それともただのせいなのかわかりません私のGPU（nvidia GTX 275）。または、pyOpenCl のオーバーヘッドが原因である場合。しかし、結果を GPU から RAM にコピーするだけでどれくらいの時間がかかるかを測定しましたが、それは行列乗算時間の約 10% にすぎません。 

#define BLOCK_SIZE 22 
__kernel void matrixMul(
      __global float* Cij, 
      __global float* Aik, 
      __global float* Bkj, 
      __const int ni, 
      __const int nj,
      __const int nk
){
//   WARRNING : interchange of  i  and  j  dimension  lower the performance >2x on my nV GT275 GPU    
int gj = get_global_id(0);    int gi = get_global_id(1); 
int bj = get_group_id(0);     int bi = get_group_id(1);  // Block index
int tj = get_local_id(0);     int ti = get_local_id(1);  // Thread index
int oj = bi*BLOCK_SIZE;       int oi = bj*BLOCK_SIZE; 
float Csub =0; 
__local float As   [BLOCK_SIZE][BLOCK_SIZE];
__local float Bs   [BLOCK_SIZE][BLOCK_SIZE];
for (int ok = 0; ok < nk; ok += BLOCK_SIZE )   {
    As[ti][tj] = Aik[ nk*(gi   ) + tj + ok ];   // A[i][k]
    Bs[ti][tj] = Bkj[ nj*(ti+ok) + gj ];        // B[k][j]
    barrier(CLK_LOCAL_MEM_FENCE);
    for (int k = 0; k < BLOCK_SIZE; ++k) Csub += As[ti][k] * Bs[k][tj];
    barrier(CLK_LOCAL_MEM_FENCE);
}
Cij[ nj * ( gi ) + gj ] = Csub;

}

注 - 奇妙な BLOCK_SIZE=22 は最大の BLOCK_SIZE であり、GPU では 512 である最大の work_group_size に適合します。このコードでは、条件 BLOCK_SIZE^2 < max work_group_size を保持する必要があります。22=int(sqrt(512))。BLOCK_SIZE=16 または 8 も試しましたが、22 では日焼けが遅くなりました。

また、単純な matrixMul (ローカル メモリを使用せずに) も試しましたが、numpy.dot() よりも 10 倍も遅くなりました。ここにコードをコピーしました http://gpgpu-computing4.blogspot.cz/2009/10/matrix-multiplication-3-opencl.html シンプルなバージョン (ローカル メモリなし) でさえ、CPU より 200 倍速く実行する必要があると彼らは言いますか? 私はそれを理解していません。

私の場合のパフォーマンスの依存関係は次のとおりです。

N =  220 numpy 3.680 [Gflops] GPU 16.428 [Gflops] speedUp 4.464 
N =  330 numpy 4.752 [Gflops] GPU 29.487 [Gflops] speedUp 6.205 
N =  440 numpy 4.914 [Gflops] GPU 37.096 [Gflops] speedUp 7.548 
N =  550 numpy 3.849 [Gflops] GPU 47.019 [Gflops] speedUp 12.217 
N =  660 numpy 5.251 [Gflops] GPU 49.999 [Gflops] speedUp 9.522 
N =  770 numpy 4.565 [Gflops] GPU 48.567 [Gflops] speedUp 10.638 
N =  880 numpy 5.452 [Gflops] GPU 44.444 [Gflops] speedUp 8.152 
N =  990 numpy 4.976 [Gflops] GPU 42.187 [Gflops] speedUp 8.478 
N = 1100 numpy 5.324 [Gflops] GPU 83.187 [Gflops] speedUp 15.625 
N = 1210 numpy 5.401 [Gflops] GPU 57.147 [Gflops] speedUp 10.581 
N = 1320 numpy 5.450 [Gflops] GPU 48.936 [Gflops] speedUp 8.979

注 - 「Gflops」数は N^3/時間として取得され、GPU からメイン メモリに結果をコピーするのに必要な時間が含まれますが、特に N>1000 の場合、この時間は合計時間の数パーセントにすぎません

多分もっと絵に描いたのは秒単位の時間です:

N =  220 numpy 0.003 [s] GPU 0.001 [s] load 0.001 [s] speedUp 5.000 
N =  330 numpy 0.008 [s] GPU 0.001 [s] load 0.001 [s] speedUp 7.683 
N =  440 numpy 0.017 [s] GPU 0.002 [s] load 0.001 [s] speedUp 7.565 
N =  550 numpy 0.043 [s] GPU 0.004 [s] load 0.001 [s] speedUp 11.957 
N =  660 numpy 0.055 [s] GPU 0.006 [s] load 0.002 [s] speedUp 9.298 
N =  770 numpy 0.100 [s] GPU 0.009 [s] load 0.003 [s] speedUp 10.638 
N =  880 numpy 0.125 [s] GPU 0.010 [s] load 0.000 [s] speedUp 12.097 
N =  990 numpy 0.195 [s] GPU 0.015 [s] load 0.000 [s] speedUp 12.581 
N = 1100 numpy 0.250 [s] GPU 0.031 [s] load 0.000 [s] speedUp 8.065 
N = 1210 numpy 0.328 [s] GPU 0.031 [s] load 0.000 [s] speedUp 10.581 
N = 1320 numpy 0.422 [s] GPU 0.047 [s] load 0.000 [s] speedUp 8.979

async_work_group_copy または read_imageui を使用してブロックをローカルメモリにコピーすることで、速度を向上できるのではないかと考えていました。しかし、200倍のスピードアップがあると言っている人と基本的に同じコードを使用しているのに、なぜそんなに大きな違いがあるのか​​ わかりません?????

4

2 に答える 2

1

優れた GPU 行列乗算は、ローカル メモリを使用するだけでなく、A、B、および/または C のブロックをレジスタに格納します (その結果、レジスタの使用率が高くなり、占有率が低くなりますが、最終的にははるかに高速になります)。これは、GPU にはローカル メモリよりも多くのレジスタがあり (128 ～ 256KB に対して NVIDIA では 48KB)、レジスタは ALU が処理できる帯域幅を提供するためです。

于 2013-10-07T14:34:22.003 に答える