OpenCLカーネル（DSPウィンドウ機能）のプロファイリングの結果について混乱しています

Question

OpenCLでウィンドウ関数カーネルを完成させました。基本的に、ウィンドウ関数は、別の数字の部分で一連の係数を1断片に適用するだけです（Wikipediaはそれをよりよく説明します）。ほとんどの場合、ウィンドウ係数フロートアレイを一定のキャッシュで詰めることができました。

Compute Profの結果は、ホストメモリ転送にデバイスとデバイスのホストが処理時間の95％以上かかることを示すことを期待していました。私のほぼすべてのケースでは、処理時間のわずか80％です。私はボードとの間の420万のフロートアレイを書いて読んでいて、一般的に100万未満のままである別のフロートアレイを書いています。

カーネルの何かは怪しげに見えますか？そもそもCPUよりもGPUでより速く実行される問題であるかどうかについての意見はありません（これについてはまだ100％ではありません）。なぜ私のgld_efficiencyとgst_efficiencyが0.1〜0.2の間でホバーする理由について少し驚いています。 G80グローバルメモリを念頭に置いてこのカーネルを作成しました。私のグローバルメモリ全体のスループットは、40GBSで大丈夫のようです。カーネルは非常にシンプルで、以下に投稿されています。

__kernel void window(__global float* inputArray, // first frame to ingest starts at 0.  Sized to nFramesToIngest*framesize samples
    __constant float* windowArray, // may already be partly filled
    int windowSize, // size of window frame, in floats
    int primitivesPerDataFrame, //amount of primitives in each frame of inputArray parameter
    int nInFramesThisCall, //each thread solves a frame, so this integer represent how many threads this kernel launches
    int isRealNumbers //0 for complex, non-zero for real 
)
{
int gid = get_global_id(0) + get_global_size(0) * get_global_id(1);

if(gid < nInFramesThisCall) //make sure we don't execute unnecessary threads
{
    if(isRealNumbers)
    {
        for(int i = 0; i < primitivesPerDataFrame; i++)
        {
            int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
            inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize];
        }
    }
    else //complex
    {
        for(int i = 0; i < primitivesPerDataFrame; i++)
        {
            int inputArrayIndex = (gid*primitivesPerDataFrame)+i;
            inputArray[inputArrayIndex] = inputArray[inputArrayIndex] * windowArray[i%windowSize/2];
        }
    }
}

}

Answer 1

ちなみに、いくつのスレッド（OPENCL用語は作業項目です）を使用していますか？大きなGPUを効率的にロードするには、少なくとも数百の何かが必要です。

合体メモリアクセスを使用したいのですが、

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

ほとんどの場合、これを可能にしません。 NvidiaのG80には、合体に関してはかなり厳しい制限があります。詳細については、「OpenCL Best Practices Guide」を参照してください。基本的に、1つのワープからの作業項目は、64または128バイトのアラインドブロックの要素に同時に特定の方法でアラインされたブロックの要素にアクセスする必要があり、負荷とストアが合体します。

またはあなたに例を与えるために：if primitivesPerDataFrame IS 16、ワープの荷重とストアは、16個の要素間隔でオフセットで行われ、効率的な合体が不可能になります。