Shoeisha Technology Media

CodeZine(コードジン)

記事種別から探す

CUDA:「超並列コンピューティング」はじめました(その2)
~高速化TipsとOpenGLでのグラフィクスへの応用

  • LINEで送る
  • このエントリーをはてなブックマークに追加
2015/03/16 14:00

 引き続きCUDAのおハナシ。前回はほんの導入編、CUDAによる超並列コンピューティングの"はじめの一歩"でした。これで完結ではあまりにつまんないし、ビデオ・カードを貸してくれたNVIDIAさんに申し訳がたちません。パフォーマンス向上のTipsとグラフィクスへの応用例を紹介します。

目次

データ転送を速くする

 前回のサンプル:1000×1000行列のかけ算の実行結果を再掲します。

補足

 簡単なサンプルとして行列積を取り上げましたけど、実アプリではこんなのわざわざ実装する必要はありません。基本的な線型代数ルーチンを集めたBLAS(Basic Linear Algebra Subprograms)のCUDA実装:cuBLASがCUDA Toolkitに含まれています。デバイス側コードはcuBLASライブラリの中にあるので、.cuを書くこともなく通常のC/C++ライブラリとしてリンクし、利用できます(cuBLASによる行列積をサンプルファイルに納めました)。

 デバイス内の処理は20GFLOPS超のナイスなスピードを叩き出しています。が、ホスト-デバイス間のデータ転送を含めると……ちょっと残念。データ転送時間が足を引っ張っています。PCI-busを越えて大量のデータをやりとりするのはそれなりの時間がかかるものです。

 ホスト-デバイス間のデータ転送はDMA(Direct Memory Access)を用い、CPUの動作とは独立してホスト側のメモリとビデオ・メモリとの間でコピーを行っています。ホスト側プログラム(プロセス)は仮想アドレス空間で動いており、プログラム内のポインタが指す値(アドレス)と実メモリ上の物理アドレスは一致しません。さらにOSはメモリの使用状況を監視していて、必要に応じてメモリの内容をディスクに保存し(ページアウト)、他の目的で使ったりします。

 DMAにしてみればページアウト領域がコピー元/先では意図しない値がコピーされることとなり、困ってしまいます。そこで物理メモリ上に特定の固定領域を用意し、コピー元からいったんそちらに転送したのちにDMAがデバイスにコピーを行います(逆方向も同様)。CUDAではページアウトしない"ピン留め"された領域(page-lockedメモリ/pinnedメモリ)を確保するAPIが用意されています。ピン留めされた領域なら、そこから直接DMA転送できるってスンポーです。

 page-lockedメモリの確保は簡単、cudaMemcpyによるホスト-デバイス間コピーの対象となるホスト側領域をcudaMallocHostで確保するだけ。

型* p = new 型[要素数];

型* p = (型*)malloc(要素数*sizeof(型));

型* p = nullptr;
cudaMallocHost(&p, 要素数*sizeof(型));

に置き換え、delete[] p;やfree(p)はcudaFreeHost(p)に改めます。

 早速実験。10メガbyteのデータをホスト-デバイス間で100往復させ、所要時間を測ります。

list01
#include "cuda_runtime.h"

#include <iostream>
#include <chrono>
#include <cassert>

using namespace std;

chrono::milliseconds memory_test(char* host, char* device, size_t size, unsigned int repeat) {
  chrono::high_resolution_clock::time_point start = chrono::high_resolution_clock::now();
  while ( --repeat ) {
    cudaMemcpy(device, host, size, cudaMemcpyHostToDevice);
    cudaMemcpy(host, device, size, cudaMemcpyDeviceToHost);
  }
  chrono::high_resolution_clock::time_point stop = chrono::high_resolution_clock::now();
  return chrono::duration_cast<chrono::milliseconds>(stop - start);
}

int main() {
    const size_t size = 1024*1024*10; // 10M-bytes
    const unsigned int repeat = 100;  // 100 times

    char* device;
    cudaMalloc(&device, size);

    chrono::milliseconds duration;

    char* host = nullptr;

  for ( int i = 0; i < 5; ++i ) {
    cout << "--- normal -------- ";
    host = new char[size];
    assert( host != nullptr );
    duration = memory_test(host, device, size, repeat);
    long long nms = duration.count();
    cout << duration.count() << " [ms]" << endl;
    delete[] host;
    host = nullptr;

    cout << "--- page-locked --- ";
    cudaError_t status = cudaMallocHost(&host, size);
    assert( status == cudaSuccess );
    duration = memory_test(host, device, size, repeat);
    long long pms = duration.count();
    cout << duration.count() << " [ms] (" << pms * 1.0f / nms << ")" << endl;
    cudaFreeHost(host);
    host = nullptr;
    cout << endl;
  }

    return 0;
}

 40%ほど時間短縮できているみたいです。この差は決して小さくないですよ。

 ホスト-デバイス間のデータ転送はその量/回数ともにできるだけ少なくなるような実装設計がパフォーマンスのキモとなりそうです。いったんホストからデバイスにデータを投げたら、極力デバイス内で処理を完結し、最終結果のみをデバイスから引き取るのがベストでしょうね。


  • LINEで送る
  • このエントリーをはてなブックマークに追加

著者プロフィール

  • επιστημη(エピステーメー)

    C++に首まで浸かったプログラマ。 Microsoft MVP, Visual C++ (2004.01~) だったり わんくま同盟でたまにセッションスピーカやったり 中国茶淹れてにわか茶人を気取ってたり、 あと Facebook とか。 著書: - STL標準講座 (監修) -...

All contents copyright © 2006-2017 Shoeisha Co., Ltd. All rights reserved. ver.1.5