cudaMemcpyの時間を計測してみる

cudaMemcpyの実行時間を計測してみました。環境条件は次の通りです。

CUDA 2.3
GeForce 9800GT

1KB転送

とりあえず1KBの転送をしてみます。

コード

#include <algorithm>
#include <cstdio>
#include <ctime>
#include <stdint.h>

#define N 1024

static inline void
print_msec(const char * s, clock_t c)
{
    double msec = (static_cast<double>(c) / CLOCKS_PER_SEC) * 1000;
    printf("%s : %f\n", s, msec);
}

int main(void)
{
    uint8_t * d_tmp = NULL;
    uint8_t input[N];
    uint8_t output[N];

    for (int i = 0; i < N; ++i) input[i] = i & 0xFF;
    std::fill(output, output + N, 0x00);

    clock_t c0 = clock();
    cudaMalloc(&d_tmp, N);
    clock_t c1 = clock();
    for(int i = 0; i < N * N; ++i) cudaMemcpy(d_tmp, input, N, cudaMemcpyHostToDevice);
    clock_t c2 = clock();
    for(int i = 0; i < N * N; ++i) cudaMemcpy(output, d_tmp, N, cudaMemcpyDeviceToHost);
    clock_t c3 = clock();
    //for (int i = 0; i < N; ++i) std::cout << static_cast<int>(output[i]) << std::endl;
    cudaFree(d_tmp);
    clock_t c4 = clock();

    print_msec("cudaMalloc", c1 - c0);
    print_msec("cudaMemcpyHostToDevice", c2 - c1);
    print_msec("cudaMemcpyDeviceToHost", c3 - c2);
    print_msec("cudaFree", c4 - c3);

    return 0;
}

実行結果

cudaMalloc : 50.000000
cudaMemcpyHostToDevice : 11070.000000
cudaMemcpyDeviceToHost : 10780.000000
cudaFree : 0.000000

1KBのデータを1024*1024回送受信してます。これを見る限り上りと下りの速度差はないようです。1KB転送の平均時間は大体10usでしょうか？心配になったのでtimeコマンドで測って見ましたが大体22sだったので計測結果の表示がは間違っている分けではなさそうです。
大体100MB/sくらいでしょうか？メインメモリではなくデバイスドライバとPCI Expressを経由して流れていく事を考えると、簡単に試した段階ではこんなもんですかね。

1MB転送

1KBという細切れで送っているせいでレイテンシが大きくでているのかとも思い、試しにデータサイズを1MBして送受信してみました。

コード

#include <algorithm>
#include <cstdio>
#include <ctime>
#include <stdint.h>

#define N (1024 * 1024)

static inline void
print_msec(const char * s, clock_t c)
{
    double msec = (static_cast<double>(c) / CLOCKS_PER_SEC) * 1000;
    printf("%s : %f\n", s, msec);
}

int main(void)
{
    uint8_t * d_tmp = NULL;
    uint8_t input[N];
    uint8_t output[N];

    for (int i = 0; i < N; ++i) input[i] = i & 0xFF;
    std::fill(output, output + N, 0x00);

    clock_t c0 = clock();
    cudaMalloc(&d_tmp, N);
    clock_t c1 = clock();
    for(int i = 0; i < 1024; ++i) cudaMemcpy(d_tmp, input, N, cudaMemcpyHostToDevice);
    clock_t c2 = clock();
    for(int i = 0; i < 1024; ++i) cudaMemcpy(output, d_tmp, N, cudaMemcpyDeviceToHost);
    clock_t c3 = clock();
    //for (int i = 0; i < N; ++i) std::cout << static_cast<int>(output[i]) << std::endl;
    cudaFree(d_tmp);
    clock_t c4 = clock();

    print_msec("cudaMalloc", c1 - c0);
    print_msec("cudaMemcpyHostToDevice", c2 - c1);
    print_msec("cudaMemcpyDeviceToHost", c3 - c2);
    print_msec("cudaFree", c4 - c3);

    return 0;
}

実行結果

cudaMalloc : 50.000000
cudaMemcpyHostToDevice : 770.000000
cudaMemcpyDeviceToHost : 760.000000
cudaFree : 0.000000

今度は1MBを1024回送受信しました。だいたい1.3GB/sでしょうか。やはり細切れのレイテンシが遅くしていたようです。メインメモリもDDR2で実測すると一桁前半GB/sくらいなので、外部機器に流すということからすると1.3GB/sかなりいい成績な気がします。

512MB転送

さらに大きくして512MBを一回で送受信してみます。

コード

#include <algorithm>
#include <cstdio>
#include <ctime>
#include <stdint.h>

#define N (512 * 1024 * 1024)

uint8_t input[N];
uint8_t output[N];

static inline void
print_msec(const char * s, clock_t c)
{
    double msec = (static_cast<double>(c) / CLOCKS_PER_SEC) * 1000;
    printf("%s : %f\n", s, msec);
}

int main(void)
{
    uint8_t * d_tmp = NULL;

    for (int i = 0; i < N; ++i) input[i] = i & 0xFF;
    std::fill(output, output + N, 0x00);
    clock_t c0 = clock();
    cudaMalloc(&d_tmp, N);
    clock_t c1 = clock();
    cudaMemcpy(d_tmp, input, N, cudaMemcpyHostToDevice);
    clock_t c2 = clock();
    cudaMemcpy(output, d_tmp, N, cudaMemcpyDeviceToHost);
    clock_t c3 = clock();
    //for (int i = 0; i < N; ++i) std::cout << static_cast<int>(output[i]) << std::endl;
    cudaFree(d_tmp);
    clock_t c4 = clock();

    print_msec("cudaMalloc", c1 - c0);
    print_msec("cudaMemcpyHostToDevice", c2 - c1);
    print_msec("cudaMemcpyDeviceToHost", c3 - c2);
    print_msec("cudaFree", c4 - c3);

    return 0;
}

実行結果

cudaMalloc : 40.000000
cudaMemcpyHostToDevice : 230.000000
cudaMemcpyDeviceToHost : 230.000000
cudaFree : 0.000000

何気なく512MBの配列2個をauto変数で取ろうしたらsegmentation faultしちゃいました。気をつけましょう。それはいいとして、512MBで230msなのでだいたい2.2GB/sでしょうか。メインメモリに対して遅いmemcpy()実行するよりも、場合によっては速いかもしれません。

まとめ

Cell/B.E.は25.6GB/sなのでそれに比べるとまだ1桁遅いわけですが、一万円くらいで買える（本来のGPUの用途の為に既に持ってる人もいる）PCのアクセラレータとして、2.2GB/sという転送速度はいい速度なんじゃないでしょうか。PCI Expressを経由するのでデータのやりとりが完全にネックになって、肝心の演算器を全然使いこなせない可能性を危惧していたのですが、これだけ速度が出るならなにか使える場面はありそうです。