6

OpenCLforAMDとNvidiaGPUのパフォーマンスの違いを評価しようとしています。行列とベクトルの乗算を実行するカーネルがあります。現在、カーネルを2つの異なるシステムで実行しています。1つはUbuntu12.04とCUDA4.0（OpenCLライブラリとヘッダーを含む）を搭載したNVidia GT525mを搭載したラップトップで、もう1つはUbuntuを搭載したAMDRadeonHD7970を搭載したデスクトップです。 12.04および最新のCatalystドライバー。

カーネルには#pragma unroll、Nvidia OpenCL実装の大幅な高速化（〜6x）を実現する2つのステートメントがあります。ただし、AMDOpenCLバージョンではスピードアップは発生しません。AMD APPカーネルアナライザーでカーネルを見ると、トリップカウントが不明なため、展開が使用されていないというエラーが表示されます。だから私の質問は、#pragma unrollAMD OpenCLで動作するのか、それとも代替手段があるのか​​（おそらく私が知らないコンパイラフラグ）です。以下にカーネルを含めました

__kernel void mvKernel(__global float* a, const __global float* x, __global float* y, int m, int n)
{
    float sum = 0.0f;
    __global float* A;
    int i;
    int j = 0;
    int indx = get_global_id(0);
    __local float xs[12000];
#pragma unroll 
    for(i = get_local_id(0); i < n; i+= get_local_size(0)) {
        xs[i] = x[i];
    } 
    barrier(CLK_LOCAL_MEM_FENCE);
    A = &a[indx];
#pragma unroll 256
    for(i = 0; i < n; i++) {
        sum += xs[i] * A[j];
        j += m;
    }
    y[indx] = sum;
}

この同じカーネルは両方の実装で正しい結果を生成しますが、＃pragma unrollコマンドはAMDに対して何もしません（コメントアウトしてチェックアウト）。

4

1 に答える 1

8

文書化されていませんが、実際にはで動作するはず#pragma unrollです。コンパイラログをチェックして、展開が適用されているかどうかを確認できますか？カーネルアナライザーがOpenCLランタイムと同じコンパイラーを使用しているかどうかはわかりませんが、確認することをお勧めします。

それ以外の場合、n256のチャンクで提供されることがわかっている場合は、256要素のブロック上に1つのループを作成し、256の固定サイズで内部に別のループを作成することで手動で展開できます。これにより、展開が簡単になる場合があります。これにより、トリップカウントが静的にわからないという問題が確実に解決されます。

ただし、計算をキャッシュするレジスタがあまりないため、ループの展開は通常、それほど大きなメリットにはならないことに注意してください。ループ展開によるレジスター圧力の増加は、レジスターの流出につながる可能性があり、これはさらに遅くなります。カーネルがAMDカード上で実際にどれだけ高速であるかを確認する必要があります。新しいNVIDIAOpenCLコンパイラも、アンロールプラグマの恩恵を受けられなくなる可能性があります。

于 2012-11-19T20:19:43.240 に答える