1

OpenCL でウィンドウ関数カーネルを完成させました。基本的に、ウィンドウ関数は、係数のセットを別の数値のセットに 1 つずつ適用するだけです (ウィキペディアの説明の方が適切です)。ほとんどの場合、ウィンドウ係数浮動小数点配列を定数キャッシュに詰め込むことができました。

Compute Prof からの結果は、ホストからデバイスへ、およびデバイスからホストへのメモリ転送に 95% 以上の処理時間がかかることを示していると予想していました。ほとんどすべてのケースで、処理時間はわずか 80% です。私は、420 ​​万の float 配列をボードとの間で読み書きし、通常は 100 万をはるかに下回る別の float 配列を書き込んでいます。

カーネル内に怪しいものはありますか? そもそもCPUよりもGPUで高速に実行する必要がある問題であるかどうかについての意見（私はまだこれについて100％ではありません）。私の gld_efficiency と gst_efficiency が 0.1 から 0.2 の間で変動する理由について、私は少し唖然としています。このカーネルは、G80 グローバル メモリの合体​​を念頭に置いて作成しました。グローバル メモリの全体的なスループットは 40 GB で問題ないようです。カーネルは非常にシンプルで、以下に掲載されています。

__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];
        }
    }
}

}

4

1 に答える 1

1

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

合体したメモリアクセスを利用したいが、次のようなオフセットのある負荷

int inputArrayIndex = (gid*primitivesPerDataFrame)+i;

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

または、例を挙げます。primitivesPerDataFrameが16の場合、ワープのロードとストアは16要素離れたオフセットで行われるため、効率的な合体は不可能です。

于 2010-11-11T16:58:05.633 に答える