CUDA by Exampleを読んでます．
Amazon.co.jp： CUDA by Example 汎用GPUプログラミング入門: Jason Sanders, Edward Kandrot, 株式会社クイープ: 本
その中で，ホスト側のメモリ確保のためにc標準のmallocと，
CUDAに用意されているcudaMallocHostの違いが説明されていますが，
その中の「ページング」という単語がわからないので調べた限りのところを以下にまとめます．
ちなみにpage-lockedﾎﾆｬﾗﾗとかの単語で，cudaのprogramming guideにも出てくるので，
割りと基本的な概念なようです．
辞書的説明
e-wordsによる説明
ページングとは｜paging - 意味/解説/説明/定義 ： IT用語辞典
ページングとは、仮想記憶（仮想メモリ）の方式の一つで、メモリ領域をページと呼ばれる一定の大きさの領域に分割し、物理的なアドレス（番地）とは別に仮想的なアドレスを割り当てて管理する方式。細切れのメモリ空間を連結して一つの連続した空間として利用したり、補助記憶装置（ハードディスクなど）上にも仮想的なメモリ領域を確保することで、物理メモリの容量を超えてメモリ空間を利用することができる。ページの大きさはOSやハードウェアによって異なるが、現代では多くのシステムが4KBのページを採用している。
各ページには物理メモリ上での所在地（物理アドレス、実アドレス）とは別に、OSによって仮想メモリ空間上でのアドレス（仮想アドレス、論理アドレス）が与えられ、プログラム（プロセス）へのメモリの割当も仮想メモリ空間上でページ単位で行われる。こうすることで物理メモリ上で細切れの領域を集めて一つの大きな仮想メモリ領域を確保することが可能で、また、プログラムの側で物理メモリを管理する必要がなくなる。
wikipediaによる説明
ページング方式 - Wikipedia
ページング方式 (Paging) とは、コンピュータのオペレーティングシステムにおいて記憶装置をページと呼ばれる小さな単位に分割して割り当てを行うアルゴリズム群である。仮想記憶のベースとなる設計の一つ。
物理メモリ空間および論理メモリ空間を、基本的に一定サイズのページと呼ばれる単位に分割して管理を行う。論理メモリから物理メモリ空間への対応づけはページテーブルと呼ばれる構造体で実現され、この構造体はオペレーティングシステム (OS) によって管理される。物理メモリ空間に対応づけられていない論理メモリを参照した時にはページフォルトという例外によってOS側の例外処理ルーチンに制御が移行し、OS側の管理によって適宜対応したページを二次記憶等から読み込み、テーブルを更新してその参照した命令の実行に戻る。
これを実現するハードウエアであるメモリ管理ユニット (MMU) の中にはトランスレーション・ルックアサイド・バッファ (Translation Lookaside Buffer:TLB) と呼ばれる一種のキャッシュがあり、ユニット内部ではこの対応表に基づいてメモリアドレスの対応づけを行っている。このテーブルから参照出来なかったときをTLBミスと呼ぶ。このときの処理はMMUの設計によって異なり、MMU内にはTLBのみを持ちTLBミスが即例外を起こし、OSがページテーブルを引いてTLBに追加することによってTLBミスを解決するアーキテクチャや、ページテーブル自体のフォーマットがOSが使えるビットを含めた形でMMUによって定義されていて、TLBミス時にMMU自身が与えられた物理アドレスにあるページテーブルを参照するアーキテクチャもある。
非常に柔らかく噛み砕くと...
もともと物理的な制約のあるメモリ空間を，OS管理によるページテーブルを使って，仮想的な論理メモリ空間に対応付けて用いることが「ページング」．
プログラムからは，この論理メモリだけを見て処理を行うことができる．
ページングしないことによる利点
ページングされないメモリは，page-locked memoryやpinned memoryなどと呼ばれる．
これらは，OSによってバッファが追い出されたり，移動したりしないので，
アプリケーションがそのメモリの物理アドレスにアクセスしても安全となる．
したがって，GPUはダイレクトメモリアクセス（DMA）を使ってホストとの間でデータコピーができる．
CUDAのドライバは，なんと，通常のmallocで確保されたページング可能なメモリからメモリコピーを行うときにも，
裏でDMAを使ってバッファをGPUに転送している．
つまり，2回分転送を行っている．
このことを考えると，page-lockedされたメモリを使うと転送が速くなることがわかる．
OS:Windows7
CPU:Intel Core i7 3.33GHz
GPU:GeForce780TiとQuadro5000
の環境で、
int型の配列を10万回、Host-Deviceのデータ転送を往復させて、その平均値をとってみた。
これをみると大体、要素数が2^12-2^14より大きければ、page-lockedなメモリ確保をしたほうがよさそう。
なので基本的にはmallocでホストメモリを確保して、
特別大きい場合にはcudaMallocHostを使ったほうがよいといえる。
他にもcudaMemcpyAsync()という非同期関数を用いる場合には、
page-lockedなメモリとデバイスメモリの間に限られる。
ただし、ページをlockすると、他のアプリやそもそもシステムの使える物理メモリが減少するので、
すべてのmallocをcudaMallocHostに置き換えるというようなことはできない。
printf("SIZE = %d\n", SIZE); int *data_host; int *data_host_pagelocked; int *data_dev; int *data_dev_pagelocked; data_host = (int*)malloc(sizeof(int)*SIZE); cudaMallocHost(&data_host_pagelocked, sizeof(int)*SIZE); cudaMalloc(&data_dev, sizeof(int)*SIZE); cudaMalloc(&data_dev_pagelocked, sizeof(int)*SIZE); /*time*/ float elapsed_time_ms = 0.0f; cudaEvent_t start,stop; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, 0); for(int i=0;i<N; i++){ cudaMemcpy(data_dev, data_host, sizeof(int)*SIZE, cudaMemcpyHostToDevice); cudaMemcpy(data_host, data_dev, sizeof(int)*SIZE, cudaMemcpyDeviceToHost); } cudaEventRecord(stop,0); cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsed_time_ms, start, stop); printf("non-locked: %8.2f total:ms, %8.2f ms\n", elapsed_time_ms, elapsed_time_ms/N); cudaEventDestroy(start); cudaEventDestroy(stop); /*time*/ elapsed_time_ms = 0.0f; cudaEventCreate(&start); cudaEventCreate(&stop); cudaEventRecord(start, 0); for(int i=0;i<N; i++){ cudaMemcpy(data_dev_pagelocked, data_host_pagelocked, sizeof(int)*SIZE, cudaMemcpyHostToDevice); cudaMemcpy(data_host_pagelocked, data_dev_pagelocked, sizeof(int)*SIZE, cudaMemcpyDeviceToHost); } cudaEventRecord(stop,0); cudaEventSynchronize(stop); cudaEventElapsedTime(&elapsed_time_ms, start, stop); printf("locked: %8.2f total:ms, %8.2f ms\n", elapsed_time_ms, elapsed_time_ms/N); // printf("time: %8.2f ms\n", elapsed_time_ms); cudaEventDestroy(start); cudaEventDestroy(stop); cudaFree(data_dev); cudaFree(data_dev_pagelocked); cudaFreeHost(data_host_pagelocked); free(data_host);
