GF100のキャッシュの性能

Motivation

GTX480に代表されるGF100およびGTX460などのGF104にはL1, L2キャッシュが搭載され、メモリアクセスに関する制限が緩和されました。何も考えずに書くとGT200系では性能が出なかった場合でも、十分な性能が発揮されることがあります。

そこで、定数がテーブルに入っていて、それを参照して計算するというコードを例に、キャッシュの性能を測定します。

普通に書く

__device__
float f(float const x) {
    return 2 * x * x;
}

__global__
void kernel(float2 const* table, float* result) {
    result += blockIdx.x * 256 + threadIdx.x;

    float sum = 0;
    for (int loop = 0; loop < 256; ++loop) {
        for (int i = 0; i < 32; ++i) {
            sum += f(table[i].x) * table[i].y;
        }
    }

    *result = sum;
}

まずは普通にCPUの様に書きます。

GT200系では、table[i]のアクセスがキャッシュされないため、メモリアクセスが毎ループ発生してしまいます。一方で、GF100系ではキャッシュによりメモリアクセスの回数軽減が見込まれる書き方です。

shared memory

__global__
void kernel(float2 const* table, float* result) {
    float2 __shared__ s[32];

    result += blockIdx.x * 256 + threadIdx.x;

    if (threadIdx.x < 32) {
        s[threadIdx.x] = table[threadIdx.x];
    }
    __syncthreads();

    float sum = 0;
    for (int loop = 0; loop < 256; ++loop) {
        for (int i = 0; i < 32; ++i) {
            sum += f(s[i].x) * s[i].y;
        }
    }

    *result = sum;
}

次はテーブルを一旦shared memoryに読み込み、それを使う書き方です。メモリアクセスは最適化されますが、本質的でない処理を記述せねばならず、開発コストは上昇してしまいます。

constant memory

float2 __constant__ table[32];

__global__
void kernel(float* result) {
    result += blockIdx.x * 256 + threadIdx.x;

    float sum = 0;
    for (int loop = 0; loop < 256; ++loop) {
        for (int i = 0; i < 32; ++i) {
            sum += f(table[i].x) * table[i].y;
        }
    }

    *result = sum;
}

次にconstant memoryを使った書き方です。参照しかしないのでこういう書き方もできます。

texture cache

texture<float2, 1, cudaReadModeElementType> tex;

__global__
void kernel(float* result) {
    result += blockIdx.x * 256 + threadIdx.x;

    float sum = 0;
    for (int loop = 0; loop < 256; ++loop) {
        for (int i = 0; i < 32; ++i) {
            float2 tmp = tex1Dfetch(tex, i);
            sum += f(tmp.x) * tmp.y;
        }
    }

    *result = sum;
}

最後にキャッシュ繋りでtexture cacheを使う書き方です。

性能比較

GTX480とGTX285での実行時間の差を示します。単位はmsです。ブロック数は32768、スレッド数は256で統一しています。

	GTX480	GTX285
cached	461.71	8704.82
shared	426.79	1327.15
constant	426.32	1519.05
texture	1635.48	2534.30

という結果になりました。2つのボードでは性能差があるので、sharedの結果を1とした相対値を以下に示します。

	GTX480	GTX285
cached	1.08	6.56
shared	1.00	1.00
constant	1.00	1.14
texture	3.83	1.91

まず、texture cacheはこういう用途に使うのは適さないようです。そして、どちらのボードでもshared memoryを使うのが最速であり、またGF100のキャッシュを頼ってもそこそこの速度が出るようです。GT200からGF100でconstant memoryの参照速度も向上しているようです。

予想できる局所性がある場合はshared memoryを使うべきでしょうが、予測が困難な場合はL1キャッシュを頼るのも手であると考えられます。とりあえずL1を使う書き方をしておいて、あとから最適化でshared memoryを使うようにする戦略もありそうです。