質問をすることでしか得られない、回答やアドバイスがある。

15分調べてもわからないことは、質問しよう!

ただいまの
回答率

89.09%

【CUDA】cudaMemcpyに時間がかかり速度が出ないときの対処法

受付中

回答 2

投稿 編集

  • 評価
  • クリップ 0
  • VIEW 407

shukrin

score 8

前提

つい先日からCUDAの勉強を始め、こちらのサイトを参考に配列同士の足し算を行うプログラムを作ってみました。並列計算による高速化を期待していましたが、あまり思ったようにいかず困っております。

発生している問題・エラーメッセージ

512*2048個の要素を持つfloat型配列2つの各要素を足し合わせて別の配列に値を格納するようなプログラムをGPU,CPUでそれぞれ実行し、その計算時間を比較してみたのですがCPUの方が速いという結果になってしまいました。
どこがボトルネックになっているのかを調べるために、GPUによる計算時間をホストからデバイスへのメモリコピー、GPUによる計算、デバイスからホストへのメモリコピーという3つの段階に分けてそれぞれにかかった時間を計測しました。以下がその結果です。

memcpy(HtD): 3314 microsec
gpu_calc: 451 microsec
memcpy(DtH): 783 microsec
cpu_calc: 1999 microsec

計算自体は一応高速化されているようですが、ホストとデバイス間のメモリコピーに時間がかかっていたのが原因でGPUの計算時間が増大していたようです。

該当のソースコード

上記サイトのソースコードを時間計測用に一部改変したものとなっています。

#include <iostream>
#include <stdio.h>
#include <algorithm>
#include <cuda_runtime.h>
#include <device_launch_parameters.h>
#include <chrono>

//N = 512×2048
int N = 1 << 20;

void vec_sum_c(float k, float c[], float a[], float b[]) {

    for (int i = 0; i < N; i++) {
        c[i] = k * a[i] + b[i];
    }
}

__global__
void vec_sum_k(float k, float* a, float* b, float* c)
{
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    c[i] = k * a[i] + b[i];
}

int main() {


    //a,b,cはホスト用、d_a,d_b,d_cはデバイス用のポインタ
    float* a, * b, * c, * d_a, * d_b, * d_c;

    //ホスト側の配列を用意
    //a = (float*)malloc(N * sizeof(float));
    //b = (float*)malloc(N * sizeof(float));
    //c = (float*)malloc(N * sizeof(float));
    cudaMallocHost(&a, N * sizeof(float));
    cudaMallocHost(&b, N * sizeof(float));
    cudaMallocHost(&c, N * sizeof(float));

    //デバイス側の配列を用意
    cudaMalloc(&d_a, N * sizeof(float));
    cudaMalloc(&d_b, N * sizeof(float));
    cudaMalloc(&d_c, N * sizeof(float));

    //a,bの配列にそれぞれ1,2を代入し、cを初期化
    for (int i = 0; i < N; i++) {
        a[i] = 1.0f;
        b[i] = 2.0f;
        c[i] = 0.0f;
    }

    //時間計測開始
    std::chrono::system_clock::time_point start, memcpyh2d, memcpyd2h, gpus, gpue, cpu;
    std::cout << "-----TimeStart-----" << std::endl;

    for (int i = 0; i < 10; i++) {

        start = std::chrono::system_clock::now();

        //ホスト側の配列の内容をデバイス側にコピー
        cudaMemcpy(d_a, a, N * sizeof(float), cudaMemcpyHostToDevice);
        cudaMemcpy(d_b, b, N * sizeof(float), cudaMemcpyHostToDevice);
        cudaMemcpy(d_c, c, N * sizeof(float), cudaMemcpyHostToDevice);

        //ホストの配列をデバイス側にコピーするのにかかった時間
        memcpyh2d = std::chrono::system_clock::now();

        //スレッドの設定
        int blocksize = 256;

        //ブロックあたりのスレッド数(blocksize)を512、
        //ブロックの総数(gridsize)をN/512用意する
        //したがって総スレッド数は blocksize × gridsize = N 個
        dim3 block(blocksize, 1, 1);
        dim3 grid(N / block.x, 1, 1);

        //GPUの計算時間計測開始
        gpus = std::chrono::system_clock::now();

        // カーネル関数の呼び出し
        vec_sum_k << <grid, block >> > (2.0f, d_a, d_b, d_c);
        cudaDeviceSynchronize();

        //GPUの計算時間計測終了
        gpue = std::chrono::system_clock::now();

        //計算結果をホストへコピー
        cudaMemcpy(c, d_c, N * sizeof(float), cudaMemcpyDeviceToHost);

        //計算結果をホストにコピー。GPUの計算終わり。
        memcpyd2h = std::chrono::system_clock::now();

        // 計測時間の表示
        std::cout << "-------GPULoop: " << i << " -------" << std::endl;
        double HtDtime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(memcpyh2d - start).count());
        std::cout << "memcpy(HtD): " << HtDtime << " microsec" << std::endl;
        double gputime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(gpue - gpus).count());
        std::cout << "gpu: " << gputime << " microsec" << std::endl;
        double DtHtime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(memcpyd2h - gpue).count());
        std::cout << "memcpy(DtH): " << DtHtime << " microsec" << std::endl;

    }

    //計算結果をホストにコピー。GPUの計算終わり。CPUの計算開始
    memcpyd2h = std::chrono::system_clock::now();

    vec_sum_c(2.0f, c, a, b);

    //CPUの計算時間計測終了
    cpu = std::chrono::system_clock::now();

    float maxError = 0.0f;    

    //計算結果の確認
    for (int i = 0; i < N; i++) maxError = std::max(maxError, abs(c[i] - 4.0f));
    std::cout << "MaxError: : " << maxError << std::endl;

    //計測時間の表示
    std::cout << "-------CPU-------" << std::endl;
    double cputime = static_cast<double>(std::chrono::duration_cast<std::chrono::microseconds>(cpu - memcpyd2h).count());
    std::cout << "cpu: " << cputime << " microsec" << std::endl;

    //メモリの開放
    free(a);
    free(b);
    free(c);

    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);

    return 0;
}

実現したいこと

メモリコピーの時間を含めてCPUより高速で動作するようなCUDAプログラムを作成するためにはどのような方法があるでしょうか?
また計算自体は高速化されているとはいえ、期待したほど(十倍~数百倍以上)ではないのも気になります。上記サイトの結果では少なくとも3000倍の高速化に成功していたようです。私のプログラムがそれほど高速化されていない理由は何でしょうか? また計算の高速化のためにはどのような工夫ができるでしょうか?

補足情報(FW/ツールのバージョンなど)

OS:Windows 10
CUDA version:10.2
GPU:nVidia GeForce GTX1060 6Gb

以下がdevice_queryを実行した結果です。

 CUDA Device Query (Runtime API) version (CUDART static linking)

Detected 1 CUDA Capable device(s)

Device 0: "GeForce GTX 1060 6GB"
  CUDA Driver Version / Runtime Version          10.2 / 10.2
  CUDA Capability Major/Minor version number:    6.1
  Total amount of global memory:                 6144 MBytes (6442450944 bytes)
  (10) Multiprocessors, (128) CUDA Cores/MP:     1280 CUDA Cores
  GPU Max Clock rate:                            1709 MHz (1.71 GHz)
  Memory Clock rate:                             4004 Mhz
  Memory Bus Width:                              192-bit
  L2 Cache Size:                                 1572864 bytes
  Maximum Texture Dimension Size (x,y,z)         1D=(131072), 2D=(131072, 65536), 3D=(16384, 16384, 16384)
  Maximum Layered 1D Texture Size, (num) layers  1D=(32768), 2048 layers
  Maximum Layered 2D Texture Size, (num) layers  2D=(32768, 32768), 2048 layers
  Total amount of constant memory:               65536 bytes
  Total amount of shared memory per block:       49152 bytes
  Total number of registers available per block: 65536
  Warp size:                                     32
  Maximum number of threads per multiprocessor:  2048
  Maximum number of threads per block:           1024
  Max dimension size of a thread block (x,y,z): (1024, 1024, 64)
  Max dimension size of a grid size    (x,y,z): (2147483647, 65535, 65535)
  Maximum memory pitch:                          2147483647 bytes
  Texture alignment:                             512 bytes
  Concurrent copy and kernel execution:          Yes with 2 copy engine(s)
  Run time limit on kernels:                     Yes
  Integrated GPU sharing Host Memory:            No
  Support host page-locked memory mapping:       Yes
  Alignment requirement for Surfaces:            Yes
  Device has ECC support:                        Disabled
  CUDA Device Driver Mode (TCC or WDDM):         WDDM (Windows Display Driver Model)
  Device supports Unified Addressing (UVA):      Yes
  Device supports Compute Preemption:            Yes
  Supports Cooperative Kernel Launch:            No
  Supports MultiDevice Co-op Kernel Launch:      No
  Device PCI Domain ID / Bus ID / location ID:   0 / 1 / 0
  Compute Mode:
     < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 10.2, CUDA Runtime Version = 10.2, NumDevs = 1
Result = PASS
  • 気になる質問をクリップする

    クリップした質問は、後からいつでもマイページで確認できます。

    またクリップした質問に回答があった際、通知やメールを受け取ることができます。

    クリップを取り消します

  • 良い質問の評価を上げる

    以下のような質問は評価を上げましょう

    • 質問内容が明確
    • 自分も答えを知りたい
    • 質問者以外のユーザにも役立つ

    評価が高い質問は、TOPページの「注目」タブのフィードに表示されやすくなります。

    質問の評価を上げたことを取り消します

  • 評価を下げられる数の上限に達しました

    評価を下げることができません

    • 1日5回まで評価を下げられます
    • 1日に1ユーザに対して2回まで評価を下げられます

    質問の評価を下げる

    teratailでは下記のような質問を「具体的に困っていることがない質問」、「サイトポリシーに違反する質問」と定義し、推奨していません。

    • プログラミングに関係のない質問
    • やってほしいことだけを記載した丸投げの質問
    • 問題・課題が含まれていない質問
    • 意図的に内容が抹消された質問
    • 過去に投稿した質問と同じ内容の質問
    • 広告と受け取られるような投稿

    評価が下がると、TOPページの「アクティブ」「注目」タブのフィードに表示されにくくなります。

    質問の評価を下げたことを取り消します

    この機能は開放されていません

    評価を下げる条件を満たしてません

    評価を下げる理由を選択してください

    詳細な説明はこちら

    上記に当てはまらず、質問内容が明確になっていない質問には「情報の追加・修正依頼」機能からコメントをしてください。

    質問の評価を下げる機能の利用条件

    この機能を利用するためには、以下の事項を行う必要があります。

回答 2

+2

CPUの計算に使っているvec_sum_c関数ですがループ内のコードが間違ってる気がします。
c[0] = k * a[i] + b[i];
ではなく、
c[i] = k * a[i] + b[i];
では?
この間違い方だと最適化によって全然ちがう処理時間になってしまいますよ?

投稿

  • 回答の評価を上げる

    以下のような回答は評価を上げましょう

    • 正しい回答
    • わかりやすい回答
    • ためになる回答

    評価が高い回答ほどページの上位に表示されます。

  • 回答の評価を下げる

    下記のような回答は推奨されていません。

    • 間違っている回答
    • 質問の回答になっていない投稿
    • スパムや攻撃的な表現を用いた投稿

    評価を下げる際はその理由を明確に伝え、適切な回答に修正してもらいましょう。

  • 2020/03/15 02:47

    ご回答ありがとうございます。おっしゃる通りでコードを書き間違えておりました。書き直してから改めて時間を計測してみました。

    memcpy(HtD): 2709 microsec
    gpu: 457 microsec
    memcpy(DtH): 810 microsec
    cpu: 2197 microsec

    これでCPUとGPUで同じ処理を行ったことになると思いますが、結果はこれまでとあまり変わらないようです。

    キャンセル

  • 2020/03/15 08:42

    じゃあ、最適化がかかっていないのですね。
    cudaでのホストとデバイスのメモリ転送ですが、ホスト側のメモリ確保にmallocではなくcudaMallocHost関数を使うといいかもしれません。
    これを使うとアロケートと同時にメモリをページロックしてくれるので多少早くなるかもしれません。
    あと、HOSTからDEVICEへの入力データの転送での
    cudaMemcpy(d_c, c, N * sizeof(float), cudaMemcpyHostToDevice);
    は必要ないと思います。

    キャンセル

  • 2020/03/15 18:27

    ご回答ありがとうございます。mallocをcudaMallocHostに置き換えて時間を測りなおしてみました。epistemeさんに最初のメモリ転送/カーネル実行は遅くなる場合があると教えていただいたので、cudaMallocHostに書き換えたうえで一連のGPU処理を10回繰り返して時間を計測しました。書き換える前の実行時間はepistemeさんへの返信に記載いたしましたので、もしよろしければそちらもご覧ください。
    cudaMemcpy(d_c, c, N * sizeof(float), cudaMemcpyHostToDevice)はd_cをcと同じ内容で初期化するために必要なのではないかと考え、残してあります。

    -----TimeStart-----
    -------GPULoop: 0 -------
    memcpy(HtD): 1344 microsec
    gpu: 300 microsec
    memcpy(DtH): 446 microsec
    -------GPULoop: 1 -------
    memcpy(HtD): 1184 microsec
    gpu: 221 microsec
    memcpy(DtH): 395 microsec
    -------GPULoop: 2 -------
    memcpy(HtD): 1219 microsec
    gpu: 220 microsec
    memcpy(DtH): 384 microsec
    -------GPULoop: 3 -------
    memcpy(HtD): 1204 microsec
    gpu: 233 microsec
    memcpy(DtH): 403 microsec
    -------GPULoop: 4 -------
    memcpy(HtD): 1240 microsec
    gpu: 274 microsec
    memcpy(DtH): 439 microsec
    -------GPULoop: 5 -------
    memcpy(HtD): 1182 microsec
    gpu: 236 microsec
    memcpy(DtH): 384 microsec
    -------GPULoop: 6 -------
    memcpy(HtD): 1173 microsec
    gpu: 220 microsec
    memcpy(DtH): 385 microsec
    -------GPULoop: 7 -------
    memcpy(HtD): 1173 microsec
    gpu: 220 microsec
    memcpy(DtH): 385 microsec
    -------GPULoop: 8 -------
    memcpy(HtD): 1206 microsec
    gpu: 221 microsec
    memcpy(DtH): 385 microsec
    -------GPULoop: 9 -------
    memcpy(HtD): 1325 microsec
    gpu: 445 microsec
    memcpy(DtH): 467 microsec
    MaxError: : 0
    -------CPU-------
    cpu: 2206 microsec

    キャンセル

+1

最初のメモリ転送/カーネル実行がかなり遅くなることはよくあります。

ホスト→デバイス / カーネル実行 / デバイス→ホスト 
を数回繰り返したとき、2回目以降のメモリ転送にも長い時間かかっていますか?

投稿

  • 回答の評価を上げる

    以下のような回答は評価を上げましょう

    • 正しい回答
    • わかりやすい回答
    • ためになる回答

    評価が高い回答ほどページの上位に表示されます。

  • 回答の評価を下げる

    下記のような回答は推奨されていません。

    • 間違っている回答
    • 質問の回答になっていない投稿
    • スパムや攻撃的な表現を用いた投稿

    評価を下げる際はその理由を明確に伝え、適切な回答に修正してもらいましょう。

  • 2020/03/15 18:19

    ご回答ありがとうございます。GPUによるホスト→デバイス / カーネル実行 / デバイス→ホスト を10回繰り返してそれぞれの時間を計測してみました。

    -----TimeStart-----
    -------GPULoop: 0 -------
    memcpy(HtD): 2438 microsec
    gpu: 450 microsec
    memcpy(DtH): 902 microsec
    -------GPULoop: 1 -------
    memcpy(HtD): 2643 microsec
    gpu: 434 microsec
    memcpy(DtH): 886 microsec
    -------GPULoop: 2 -------
    memcpy(HtD): 2342 microsec
    gpu: 425 microsec
    memcpy(DtH): 1176 microsec
    -------GPULoop: 3 -------
    memcpy(HtD): 2184 microsec
    gpu: 386 microsec
    memcpy(DtH): 777 microsec
    -------GPULoop: 4 -------
    memcpy(HtD): 2136 microsec
    gpu: 383 microsec
    memcpy(DtH): 777 microsec
    -------GPULoop: 5 -------
    memcpy(HtD): 2094 microsec
    gpu: 382 microsec
    memcpy(DtH): 775 microsec
    -------GPULoop: 6 -------
    memcpy(HtD): 2428 microsec
    gpu: 427 microsec
    memcpy(DtH): 881 microsec
    -------GPULoop: 7 -------
    memcpy(HtD): 2414 microsec
    gpu: 414 microsec
    memcpy(DtH): 783 microsec
    -------GPULoop: 8 -------
    memcpy(HtD): 2190 microsec
    gpu: 381 microsec
    memcpy(DtH): 785 microsec
    -------GPULoop: 9 -------
    memcpy(HtD): 2386 microsec
    gpu: 471 microsec
    memcpy(DtH): 785 microsec
    MaxError: : 0
    -------CPU-------
    cpu: 2220 microsec

    確かに数回繰り返した後の方が最初よりも少し速くなっているようです。TaroToyotomiさんにいただいたアドバイスをもとにmallocをcudaMallocHostに置き換えるとよりはっきりと効果が表れたのでよろしければそちらもご覧ください。

    キャンセル

15分調べてもわからないことは、teratailで質問しよう!

  • ただいまの回答率 89.09%
  • 質問をまとめることで、思考を整理して素早く解決
  • テンプレート機能で、簡単に質問をまとめられる

同じタグがついた質問を見る