2

カーネルのパフォーマンスに影響を与える根本的な要因を理解できないようです。2 つの単純なカーネルを実装しました。1 つは 2 つの画像を読み込んでピクセルごとに加算するカーネル、もう 1 つは 2 つの画像を読み込んでビット単位のAND 演算を行うカーネルです。ここで、カーネルが 1、3、および 4 チャネルのイメージだけでなく、8 ビットおよび 32 ビットのイメージを取得できるように、それらをテンプレート化しました。

そのため、最初は両方のカーネルでグローバル メモリをuchar3およびなどfloat3と一緒にロードしていましたuchar4。しかし、合体のためにトリプルを使用することについてあまり確信が持てなかったので、プロファイリングを実行することにしました。操作はチャネル番号に依存しないため、実際ucharの画像ではなく、幅が 3 倍の 1 チャネルの画像であるかのように画像を読み取ることができると考えましuchar3た。

実際、uchar3グローバル ロードは、ロードよりもはるかに遅くなりましたuchar。私の努力は報われました。しかし、残念ながら、これは算術カーネルでのみ発生しました。ビットごとの AND 演算は、正反対の結果を示しました。

uintこれで、画像データをs ではなくuchars としてロードするだけで、ビットごとの操作が可能になり、合体を完全に処理できることがわかりました。しかし、何が起こっているのかを学び、理解したいだけだと仮定しましょう。

float3s やs などのことは忘れましょうfloat4。私の問題はuchar、カーネルのバージョンにあります。では、簡単に言えば、なぜ負荷はuchar負荷よりも速い場合uchar3とそうでない場合があるのでしょうか?

GTX 470、コンピューティング機能 2.0 を使用しています。

PS。CUDA プログラミング ガイドによると、論理演算と加算演算のスループットは同じです。uchar(私のカーネルは実際には最初にs をs に変換するuint必要がありますが、それは両方のカーネルで行われているはずです。) したがって、実行の長さは、私が収集したものとほぼ同じになるはずです。

算術加算カーネル (ucharバージョン):

__global__ void add_8uc1(uchar* inputOne, uchar* inputTwo, uchar* output, unsigned int width, unsigned int height, unsigned int widthStep)
{
    const int xCoordinateBase = blockIdx.x * IMAGE_X * IMAGE_MULTIPLIER + threadIdx.x;
    const int yCoordinate = blockIdx.y * IMAGE_Y + threadIdx.y;

    if (yCoordinate >= height)
        return;

#pragma unroll IMAGE_MULTIPLIER
    for (int i = 0; i < IMAGE_MULTIPLIER && xCoordinateBase + i * IMAGE_X < width; ++i)
    {
        //  Load memory.
        uchar* inputElementOne = (inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x));
        uchar* inputElementTwo = (inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x));

        //  Write output.
        *(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)) = inputElementOne[0] + inputElementTwo[0];
    }
}

ビットごとの AND カーネル:

__global__ void and_8uc1(uchar* inputOne, uchar* inputTwo, uchar* output, unsigned int width, unsigned int height, unsigned int widthStep)
{
    const int xCoordinateBase = blockIdx.x * IMAGE_X * IMAGE_MULTIPLIER + threadIdx.x;
    const int yCoordinate = blockIdx.y * IMAGE_Y + threadIdx.y;

    if (yCoordinate >= height)
        return;

#pragma unroll IMAGE_MULTIPLIER
    for (int i = 0; i < IMAGE_MULTIPLIER && xCoordinateBase + i * IMAGE_X < width; ++i)
    {
        //  Load memory.
        uchar* inputElementOne = (inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x));
        uchar* inputElementTwo = (inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x));

        //  Write output.
        *(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x)) = inputElementOne[0] & inputElementTwo[0];
    }
}

uchar3ロード/ストア行が次のようになったことを除いて、バージョンは同じです。

        //  Load memory.
    uchar3 inputElementOne = *reinterpret_cast<uchar3*>(inputOne + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3);
    uchar3 inputElementTwo = *reinterpret_cast<uchar3*>(inputTwo + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3);

    //  Write output.
    *reinterpret_cast<uchar3*>(output + yCoordinate * widthStep + (xCoordinateBase + i * IMAGE_X + threadIdx.x) * 3) 
        = make_uchar3(inputElementOne.x + inputElementTwo.x, inputElementOne.y + inputElementTwo.y, inputElementOne.z + inputElementTwo.z);

AND カーネルについても同様です。（正直なところ、カーネルを正確に覚えているかどうかはわかりません...明日確認します）。

4

1 に答える 1

1

uchar3SM の命令セットには 24 ビットのロードがないため、ロードはコンパイラによって個別のロードに分割されます。そのため、それらが合体することはありません。ある程度、キャッシュはこれを軽減します。

ただし、正確な実行構成によっては、スレッドごとに約 10.7 バイトのキャッシュしかない場合があります (カーネルが単純であるため、多くのスレッドが 1 つの SM で同時に実行できるため、例はおそらくその値に近くなります)。キャッシュは完全に関連付けられていないため、スレッドごとに使用可能なバイト数は、スラッシングが発生する前にはるかに少なくなる可能性があります。正確にそれがいつ発生するかは、命令の正確なスケジューリングを含む多くの要因に依存します。これは、文書化されたスループットが同じ命令であっても異なる場合があります。

cuobjdump -sass両方のバージョンの実行可能ファイルの出力を比較して、コンパイラによる静的スケジューリングが同じかどうかを確認できます。ただし、実行時の動的スケジューリングがどのように機能するかは、基本的に観察できません。

お気づきのように、画像のすべてのチャネルは同じ方法で処理されるため、スレッド間でチャネルをどのように分配するかは問題ではありません。あなたが持っている最良のオプションは、またはuchar4の代わりに使用することです。これにより、（画像の適切な配置を前提として）キャッシュとは無関係に結合されたアクセスが得られます。これにより、実行時間が短縮され、一貫性が向上します。uchar3uchar

于 2012-11-29T01:59:37.040 に答える