0

次のCUDAカーネルをプロファイリングしています

__global__ void fftshift_2D(double2 *data, int N1, int N2)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;

    if (i < N1 && j < N2) {
        double a = pow(-1.0, (i+j)&1);

        data[j*blockDim.x*gridDim.x+i].x *= a;
        data[j*blockDim.x*gridDim.x+i].y *= a;
    }
 }

これは基本的に、2D倍精度複素数データ行列にスカラー倍精度変数を乗算します。

ご覧のとおり、私は統合されたグローバルメモリアクセスを実行しています。グローバルメモリの負荷とストアの効率を検査することにより、NVIDIAVisualProfilerでこれを確認したいと思います。驚いたことに、そのような効率は両方とも正確に50％であり、合体したメモリアクセスで期待される100％からはほど遠いことがわかりました。これは、複素数の実数部と虚数部のインターレースストレージに関連していますか？もしそうなら、100％の効率を回復するために私が利用できるトリックはありますか？

前もって感謝します。

追加情報

BLOCK_SIZE_x=16
BLOCK_SIZE_y=16

dim3 dimBlock2(BLOCK_SIZE_x,BLOCK_SIZE_y);
dim3 dimGrid2(N2/BLOCK_SIZE_x + (N2%BLOCK_SIZE_x == 0 ? 0:1),N1/BLOCK_SIZE_y + (N1%BLOCK_SIZE_y == 0 ? 0:1));

N1とN2は任意の偶数にすることができます。

カードはNVIDIAGT540Mです。

4

2 に答える 2

5

さまざまなメモリ アクセス パターンの効率性に関するこの NVIDIA ブログ投稿をご覧ください。Strided メモリ アクセスの問題が発生しています。

各コンポーネントは独立して使用されるため、代わりに配列をdouble2プレーンな通常の配列として扱うことができます( Robert Crovella が提案したように)。double

__global__ void fftshift_2D(double *data, int N1, int N2)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;

    if (i < N1 * 2 && j < N2) {
        double a = pow(-1.0, (i / 2 + j)&1);
        data[j*blockDim.x*gridDim.x+i] *= a;
    }
}

ただし、単一のスレッドで x と y の両方のコンポーネントにアクセスする必要がある場合は、次を試してください。

2 つの別々の配列を使用します。x 成分を持つもの y 成分を持つもの。そのように：

__global__ void fftshift_2D(double *dataX, double *dataY, int N1, int N2)
{
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;

    if (i < N1 && j < N2) {
        double a = pow(-1.0, (i+j)&1);

        dataX[j*blockDim.x*gridDim.x+i] *= a;
        dataY[j*blockDim.x*gridDim.x+i] *= a;
    }
}

または、データ レイアウトをそのままにして、ストライドなしで共有メモリにロードし、共有メモリから再シャッフルします。それは多かれ少なかれそのように見えます：

__global__ void fftshift_2D(double2 *data, int N1, int N2)
{
    __shared__ double buff[BLOCK_SIZE*2];
    double2 *buff2 = (double2 *) buff;
    int i = threadIdx.x + blockDim.x * blockIdx.x;
    int j = threadIdx.y + blockDim.y * blockIdx.y;
    double ptr = (double *) &data[j*blockDim.x*gridDim.x + blockDim.x * blockIdx.x];

    // TODO add guarding with N1 & N2
    buff[threadIdx.x] = ptr[threadIdx.x];
    buff[blockDim.x + threadIdx.x] = ptr[blockDim.x + threadIdx.x];
    __syncthreads();

    double a = pow(-1.0, (i+j)&1);
    buff2[threadIdx.x].x *= a 
    buff2[threadIdx.x].y *= a 

    __syncthreads();
    ptr[threadIdx.x] = buff[threadIdx.x];
    ptr[blockDim.x + threadIdx.x] = buff[blockDim.x + threadIdx.x];
}
于 2013-01-10T08:36:00.237 に答える
4

はい、構造体のデータ格納形式の配列があり、次の行で他のすべての要素のみを参照しているためです。

    data[j*blockDim.x*gridDim.x+i].x *= a;

その結果として発生するグローバル ロードとグローバル ストアの使用率はそれぞれ 50% しかありません。次の行で代替要素を参照しているため、ここではキャッシュが役立つはずです。しかし、ロード/ストアの効率はまだ 50% です。

再キャストするいくつかの方法を使用して、これを回避できると思います（この特定の例の場合）*data：

double *mydata = (double *)data;
...
mydata[2*(j*blockDim.x*gridDim.x)+i] *= a;

同じカバレッジを取得する方法を正確に示すつもりはないことに注意してください。アイデアを説明するだけです。上記のコードはほぼ必要なものですが、乗算するすべての要素が正しく処理されるようにコードを微調整する必要があります。

于 2013-01-09T22:18:52.727 に答える