編集：時間の経過に伴う成果は、この質問の最後にリストされています（まだ〜1Tflops/s）。
私は、C++ DLL から opencl(gpu) を使用して C# 用のある種の数学ライブラリを作成しており、単精度正方行列 - 行列乗法 (学習目的および後でニューラル ネットワーク プログラムでの再利用の可能性のため) でいくつかの最適化を既に行っています。以下のカーネル コードは、v1 1D 配列を matrix1(1024x1024) の行として取得し、v2 1D 配列を matrix2 の列として取得し ((1024x1024)転置最適化)、結果を v3 1D 配列にマトリックス 3 の行として配置します (1024x1024)。
今のところ、1024x1024 正方行列 - 行列乗算のカーネル実行時間は、HD7870 で 3.6 ms です。
行われた最適化:
- 2 番目の行列の転置 (改善された時間)
- 32x32 サブマトリックスを使用してローカル メモリで計算する (私の HD7870 では最大ワークグループ サイズが 256 であり、gpu は何らかの理由で 24kB を超えるローカルを受け入れないため、4x 16x16 ですが、オンライン ソースは 64kB と言っていますか?)(とにかく、かなりの差で時間を改善しました)
- 結果をローカルおよびグローバルに書き込む前に、プライベート変数を使用してデータを再利用することを増やしました。(時間の改善)
- 最も内側のループでのローカル 2D 配列への列優先アクセス。（タイムアップ）
- パッチごとに 2 つのアキュムレータ レジスタへの加算を共有します。(時間の改善と数値安定性の低下)
- 最も内側のループをループ展開しても時間は改善されませんでした (4 回目の展開後にさらに悪化しました) (したがって、整数 alu を緩和する必要があります)
質問:すべてのローカル (lds) バンク競合の排除やメモリ レイテンシを隠すための命令の並べ替えなどの最適化を完了できませんでした。この数学関数のパフォーマンスを向上させるにはどうすればよいですか?
このカーネルは確かにローカルメモリの帯域幅 (競合) に制限されており、乗算に 3.2 ミリ秒かかります =
(1024*1024*1024 * (1 sum + 1 mult =2) / 0.0036 seconds )= 596x10^9 Flops per second(596 GFlops) コンピューティング ユニットごとのローカル メモリが多いか、コアが多いか、またはその両方があるためですか?
(1024*1024*1024*(2 float 読み取り)*(4 バイト/float) /0.0036 秒)=2386x10^9 バイト/秒 しかし、このカーネルは 8 個の float を読み取り、それらを 16 回使用し、データの再利用は 2 です。フロートごと。
2386x10^9 バイト / 再利用 (2) = 1193 GB/秒
HD7870 の理論上の最大値は次のとおりです。ここでは、付録 D
計算能力 = 毎秒 2560 ギガ浮動小数点演算、LDS 帯域幅 = 2560 GB/秒、レジスタ アクセス帯域幅 = 15360 GB/秒
カーネルは次のとおりです。
__kernel void squareGpuMatrixMul(__global float * v1, __global float * v2, __global float * v3)
{
int localRow = get_local_id(0);
int localCol = get_local_id(1);
int selectRowFromA = get_group_id(0)*32;
int selectColFromB = get_group_id(1)*32;
int lid= localCol*16+localRow;
__local float Lcache1[ 16][ 16];
__local float Lcache2[ 16][ 16];
__local float Lcache3[ 16][ 16];
__local float Lcache1a[ 16][ 16];
__local float Lcache2a[ 16][ 16];
__local float Lcache3a[ 16][ 16];
__local float Lcache1b[ 16][ 16];
__local float Lcache2b[ 16][ 16];
__local float Lcache3b[ 16][ 16];
__local float Lcache1c[ 16][ 16];
__local float Lcache2c[ 16][ 16];
__local float Lcache3c[ 16][ 16];
float tmp0=0.0f;
float tmp1=0.0f;
float tmp2=0.0f;
float tmp3=0.0f;
float tmp4=0.0f;
float tmp5=0.0f;
float tmp6=0.0f;
float tmp7=0.0f;
float sumPatch=0.0f;
float sumPatcha=0.0f;
float sumPatchb=0.0f;
float sumPatchc=0.0f;
float sumPatch2=0.0f;
float sumPatcha2=0.0f;
float sumPatchb2=0.0f;
float sumPatchc2=0.0f;
barrier(CLK_LOCAL_MEM_FENCE);
Lcache3[localRow][localCol]=0.0f;
Lcache3a[localRow][localCol]=0.0f;
Lcache3b[localRow][localCol]=0.0f;
Lcache3c[localRow][localCol]=0.0f;
barrier(CLK_LOCAL_MEM_FENCE);
for(int i=0;i<1024;i+=32) // this is A's row and B's column parsed by sub-matrices
{
barrier(CLK_LOCAL_MEM_FENCE);
Lcache1[localCol][localRow]=v1[selectRowFromA*1024+i+localCol+localRow*1024];
Lcache2[localRow][localCol]=v2[selectColFromB*1024+i+localRow+localCol*1024];
Lcache1a[localCol][localRow]=v1[selectRowFromA*1024+i+localCol+localRow*1024+ 16];
Lcache2a[localRow][localCol]=v2[selectColFromB*1024+i+localRow+localCol*1024+ 16];
Lcache1b[localCol][localRow]=v1[selectRowFromA*1024+i+localCol+localRow*1024+16384];
Lcache2b[localRow][localCol]=v2[selectColFromB*1024+i+localRow+localCol*1024+16384];
Lcache1c[localCol][localRow]=v1[selectRowFromA*1024+i+localCol+localRow*1024+ 16+16384];
Lcache2c[localRow][localCol]=v2[selectColFromB*1024+i+localRow+localCol*1024+ 16+16384];
barrier(CLK_LOCAL_MEM_FENCE);
sumPatch=0.0f;
sumPatcha=0.0f;
sumPatchb=0.0f;
sumPatchc=0.0f;
sumPatch2=0.0f;
sumPatcha2=0.0f;
sumPatchb2=0.0f;
sumPatchc2=0.0f;
for(int kk=0;kk< 16;kk++) //this is sub-matrix multiplication
{
read_mem_fence(CLK_LOCAL_MEM_FENCE);
tmp0=Lcache1[kk][localRow]; // row-major
tmp1=Lcache1a[kk][localRow]; // accesses
tmp2=Lcache1b[kk][localRow]; //to local memory
tmp3=Lcache1c[kk][localRow];
tmp4=Lcache2[kk][localCol];
tmp5=Lcache2a[kk][localCol];
tmp6=Lcache2b[kk][localCol];
tmp7=Lcache2c[kk][localCol];
read_mem_fence(CLK_LOCAL_MEM_FENCE);
sumPatch+=tmp0*tmp4;
sumPatcha+=tmp0*tmp6;
sumPatchb+=tmp2*tmp4;
sumPatchc+=tmp2*tmp6;
sumPatch2+=tmp1*tmp5;
sumPatcha2+=tmp1*tmp7;
sumPatchb2+=tmp3*tmp5;
sumPatchc2+=tmp3*tmp7;
}
Lcache3[localRow][localCol]+=sumPatch+sumPatch2;
Lcache3a[localRow][localCol]+=sumPatcha+sumPatcha2;
Lcache3b[localRow][localCol]+=sumPatchb+sumPatchb2;
Lcache3c[localRow][localCol]+=sumPatchc+sumPatchc2;
}
barrier(CLK_LOCAL_MEM_FENCE);
v3[selectRowFromA*1024+selectColFromB+localCol+localRow*1024]=Lcache3[localRow][localCol];
v3[selectRowFromA*1024+selectColFromB+localCol+localRow*1024+ 16]=Lcache3a[localRow][localCol];
v3[selectRowFromA*1024+selectColFromB+localCol+localRow*1024+16384]=Lcache3b[localRow][localCol];
v3[selectRowFromA*1024+selectColFromB+localCol+localRow*1024+ 16+16384]=Lcache3c[localRow][localCol];
barrier(CLK_LOCAL_MEM_FENCE);
}
以下は、バンクの競合を排除しようとしたものですが、カーネルの実行時間は約 20 % 増加しました。
for(int kk=0;kk< 16;kk++)
{
int nc=(kk+lid)&15;//different for all local threads
//but does not exceed 0-15 range
//summation order is not important
//0.+1.+...15. or 14.+15.+0.+..13.
//gives correct answer
read_mem_fence(CLK_LOCAL_MEM_FENCE);
tmp0=Lcache1[nc][localRow];
tmp1=Lcache1a[nc][localRow];
tmp2=Lcache1b[nc][localRow];
tmp3=Lcache1c[nc][localRow];
tmp4=Lcache2[nc][localCol];
tmp5=Lcache2a[nc][localCol];
tmp6=Lcache2b[nc][localCol];
tmp7=Lcache2c[nc][localCol];
read_mem_fence(CLK_LOCAL_MEM_FENCE);
sumPatch+=tmp0*tmp4;
sumPatcha+=tmp0*tmp6;
sumPatchb+=tmp2*tmp4;
sumPatchc+=tmp2*tmp6;
sumPatch2+=tmp1*tmp5;
sumPatcha2+=tmp1*tmp7;
sumPatchb2+=tmp3*tmp5;
sumPatchc2+=tmp3*tmp7;
}
これは新しいgpuの放送技術でしょうか？また、16 要素の合計は、16 バンクのみが使用されることを意味しますか? デバイスには、ローカル アクセス用に 32 のバンクがあります。
これが私がメモリレイテンシを隠そうとしたものです:
for(int kk=0;kk< 16;kk++)
{
int nc=(kk+lid)&15;//different for all local threads
//but does not exceed 0-15 range
//summation order is not important
//0.+1.+...15. or 14.+15.+0.+..13.
//gives correct answer
read_mem_fence(CLK_LOCAL_MEM_FENCE);
tmp0=Lcache1[nc][localRow];
tmp4=Lcache2[nc][localCol];
sumPatch+=tmp0*tmp4;
tmp6=Lcache2b[nc][localCol];
sumPatcha+=tmp0*tmp6;
tmp1=Lcache1a[nc][localRow];
tmp7=Lcache2c[nc][localCol];
sumPatcha2+=tmp1*tmp7;
tmp5=Lcache2a[nc][localCol];
sumPatch2+=tmp1*tmp5;
tmp2=Lcache1b[nc][localRow];
sumPatchb+=tmp2*tmp4;
sumPatchc+=tmp2*tmp6;
tmp3=Lcache1c[nc][localRow];
sumPatchb2+=tmp3*tmp5;
sumPatchc2+=tmp3*tmp7;
read_mem_fence(CLK_LOCAL_MEM_FENCE);//this lines' position does not change time
}
しかし、これは exec を増減しませんでした。時間。
カーネル時間を改善するにはどうすればよいですか? 実行可能ですか？
デバイス: HD7870 @ 1000MHz/1200MHz ホスト: FX8150@4GHz ヘッダー、Khronos のサイトからの LIB ファイル、AMD のドライバーからの opencl.dll。
時間のサンプリングは次のように行われます: カーネルを 100 回循環させ、Stopwatchstart() および stop() メソッドからの合計時間を 100.0 で割ります。実行のみで、配列のコピーは含まれません。
すべての結果は、ランダム行列の同じ入力を使用した単純な 3 ネスト ループ バージョンと比較されます (結果は m(ij)+/-delta の範囲内で、delta は 0.001f です)。
ここでのカーネルは、より一般化されたものの単純化されたバージョンです (さまざまなマトリックスとパッチ サイズ用)。
このバージョンのカーネル パラメータ: グローバル= 512,512 ローカル=16,16、リファレンス=0,0
8320x8320 マトリックスの場合 --->Global=4160,4160, Local=16,16, ref=0,0 time = 1.87Seconds
編集:ローカルの Lcache3 をプライベート バージョンに置き換えると、DarkZeros の提案により、1024x1024 の時間が 2.7 ミリ秒に改善されました。これは毎秒 795 GFlops です。これは占有率が良いからでしょう。
Edit2:ローカルでの使用量が少ないため、48x48 (9 x 16x16) パッチを使用する可能性が開かれ、1056x1056 の乗算が 2.4 ms -->981 Gflops/s になりました。8208x8208 は、1150 GFlops を超える 961ms で実行されます。