GPUのメモリ階層の詳細（様々なメモリの利用）

第5回 GPUのメモリ階層の詳細
（様々なメモリの利用）
長岡技術科学大学
電気電子情報工学専攻
出川智啓
今回の内容

コンスタントメモリ

ページロックホストメモリ

ゼロコピーホストメモリ

ベクトル和による性能比較
392
先端GPGPUシミュレーション工学特論
2015/05/14
メモリの種類
オフチップメモリ（GPUのチップ外部のメモリ）



低速アクセス，大容量
CPUから直接アクセス可能

ローカルメモリだけはアクセス不可
グローバルメモリ
ローカルメモリ
テクスチャメモリ
コンスタントメモリ
容量
大
小
大
小
速度
低速
低速
高速*
高速*
読み書き可
読み書き可
読み込み可
読み込み可
全てのスレッドが同
じメモリにアクセス
各スレッドが異なる
メモリにアクセス
全てのスレッドが同
じメモリにアクセス
全てのスレッドが同
じメモリにアクセス
読み書き可
読み書き不可
書き込み可
書き込み可
GPUからの
読み書き
CPUからの
アクセス
393
先端GPGPUシミュレーション工学特論
*キャッシュが効く場合
2015/05/14
コンスタントメモリ


GPU全体で同じメモリに
アクセス
コンスタントキャッシュを
利用することで，効率的
な読み込みが可能

キャッシュはオンチップ
GPU
Chip
SM
SM
共有
メモリ
L1キャッ
シュ
レジ
スタ
レジ
スタ
レジ
スタ
レジ
スタ
CUDA CUDA CUDA CUDA Core Core Core Core
L1キャッ
シュ
レジ
スタ
レジ
スタ
GPU全体で64kB
レジ
スタ
ローカルローカル
メモリ
メモリ
･･･
ローカルローカル
メモリ
メモリ
･･･
コンスタントメモリ
ホスト
メモリ
テクスチャメモリ
グローバルメモリ
394
レジ
スタ
CUDA CUDA CUDA CUDA Core Core Core Core
L2キャッシュ

共有
メモリ
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリの利用


修飾子 __constant__ を付けて宣言
メモリは読込専用



CPUからは変更可能
専用のメモリ転送命令でコピー
cudaMemcpyToSymbol





395
CPU上のメモリをコンスタントメモリにコピーする
cudaMemcpyToSymbol(転送先変数名, 転送元アドレス,
バイト数, オフセット, 方向);
オフセット，方向は無くてもよい
オフセットを省略すると0が使われる
方向を省略するとcudaMemcpyHostToDeviceが使われる
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリの利用

cudaError_t cudaMemcpyToSymbol(
const char * const void * size_t
size_t
enum cudaMemcpyKind
symbol,
src,
count,
offset = 0,
kind = cudaMemcpyHostToDevice
)
 Parameters





symbol
src
count offset kind ‐
‐
‐
‐
‐
Symbol destination on device
Source memory address
Size in bytes to copy
Offset from start of symbol in bytes
Type of transfer
http://docs.nvidia.com/cuda/cuda‐runtime‐api/
396
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリの宣言

サイズは静的に決定




__constant__ 型変数名;
__constant__ 型変数名[要素数];
 配列としても宣言可能
 要素数はコンパイル時に確定している必要がある
cudaMalloc()やcudaFree()は不要
グローバル変数として宣言し，複数のカーネルから
アクセス


397
読込専用のメモリならではの使い方
書込可能なメモリでは厳禁
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリへのアクセス
（Tesla世代）

コンスタントメモリへ高速にアクセスできる要因
1. ブロードキャストによるデータの分配


16スレッド（Half Warp）単位でアクセスし，1回の読込を
他のスレッドにブロードキャストできる
16スレッドが同じアドレスにアクセスすると最も効率がよい
2. コンスタントメモリキャッシュ


398
SMごとに存在する独自のオンチップキャッシュ
他のHalf Warpがキャッシュされたデータへアクセスして
も，コンスタントメモリからの読込が発生しない
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリへのアクセス
（Tesla世代）

オフチップメモリ（DRAM）からの読込量を抑制

オフチップメモリ（DRAM）からの読込による実行速度
低下を回避

コンスタントメモリへのアクセスの制約

399
Half Warp内のスレッド全てが異なるコンスタントメモリの
アドレスを参照すると，読込が逐次化
 読込命令の処理に最悪で16倍の時間を要する
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリへのアクセス
（Tesla世代）

Half Warp内のスレッド全てが同じ
メモリアドレスにアクセス

1スレッドの読込をブロードキャストによっ
て残りのスレッドが共有
メモリアドレス
0
0
1
2
3
4
5
6

他のHalf Warpも同じメモリアドレス
にアクセス

データがキャッシュされているため，
キャッシュから高速に読み出し
7
8
9
10
11
12
13
14
15
Half Warp内でのスレッドID
400
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリへのアクセス
（Tesla世代）

Half Warp内のスレッド全てが異
なるメモリアドレスにアクセス


読込が逐次化
読込処理の時間は，Half Warpがアク
セスするコンスタントメモリアドレスの数
に比例
 最悪で処理に16倍の時間がかかる

401
おそらくグローバルメモリアクセスよりも遅
くなる
先端GPGPUシミュレーション工学特論
0
0
1
1
2
2
3
3
4
4
5
5
6
6
7
7
8
8
9
9
10
10
11
11
12
12
13
13
14
14
15
15
2015/05/14
コンスタントメモリへのアクセス
（Fermi世代）

コンスタントメモリへ高速にアクセスできる要因
1. ブロードキャストによるデータの分配


32スレッド（Warp）単位でアクセスし，1回の読込を他のス
レッドにブロードキャストできる
32スレッドが同じアドレスにアクセスすると最も効率がよい
2. コンスタントメモリキャッシュ


402
SMごとに存在する独自のオンチップキャッシュ
他のWarpがキャッシュされたデータへアクセスしても，コン
スタントメモリからの読込が発生しない
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリへのアクセス
（Fermi世代）

オフチップメモリ（DRAM）からの読込量を抑制

オフチップメモリ（DRAM）からの読込による実行速度
低下を回避

コンスタントメモリへのアクセスの制約

403
Warp内のスレッド全てが異なるコンスタントメモリのアドレ
スを参照すると，読込が逐次化
 読込命令の処理に最悪で32倍の時間を要する
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリへのアクセス
（Fermi世代）

Warp内のスレッド全てが同じコンスタントメモリアドレ
スにアクセス


1スレッドの読込をブロードキャストによって残りのスレッド
が共有
他のWarpも同じメモリアドレスにアクセス

データがキャッシュされているため，キャッシュから高速に
読み出し
0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31
0
404
16
32
48
64
80
96
先端GPGPUシミュレーション工学特論
11
2
2015/05/14
コンスタントメモリへのアクセス
（Fermi世代）

Warp内のスレッド全てが異なるコンスタントメモリアド
レスにアクセス


読込が逐次化
読込処理の時間は，Warpがアクセスするコンスタントメモリ
アドレスの数に比例
 最悪で処理に32倍の時間がかかる

おそらくグローバルメモリアクセスよりも遅くなる
0 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24 25 26 27 28 29 30 31
0
405
16
32
48
64
80
96
先端GPGPUシミュレーション工学特論
11
2
2015/05/14
コンスタントメモリ利用の例


ベクトル和
ベクトルAとBの値が全て同じ

コンスタントメモリにデータを一つ置き，全スレッドが参照
・・・
a[i]
＋
＋
＋
＋
a
＋
＋
b[i]
・・・
b
c[i]
・・・
c[i]
406
先端GPGPUシミュレーション工学特論
・・・
2015/05/14
GPUプログラム（グローバルメモリ利用）
#define N (8*1024) //64kBに収める
#define Nbytes (N*sizeof(float))
#define NT (256)
#define NB (N/NT)
__global__ void init(float *a,
float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
}
__global__ void add(float *a,
float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
c[i] = a[i] + b[i];
}
407
int main(void){
float *a,*b,*c;
cudaMalloc((void **)&a, Nbytes);
cudaMalloc((void **)&b, Nbytes);
cudaMalloc((void **)&c, Nbytes);
init<<< NB, NT>>>(a,b,c);
add<<< NB, NT>>>(a,b,c);
cudaFree(a);
cudaFree(b);
cudaFree(c);
return 0;
}
vectoradd.cu
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリ（単純な置き換え）
int i;
#define N (8*1024) //64kBに収める
#define Nbytes (N*sizeof(float))
#define NT (256)
#define NB (N/NT)
__constant__ float a[N],b[N];
__global__ void init(float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
c[i] = 0.0f;
}
__global__ void add(float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
c[i] = a[i] + b[i];
}
int main(void){
float *c;
float *host_a,*host_b;
408
host_a=(float *)malloc(Nbytes);
host_b=(float *)malloc(Nbytes);
cudaMalloc((void **)&c,Nbytes);
for(i=0;i<N;i++){
host_a[i] = 1.0f;
host_b[i] = 2.0f;
}
cudaMemcpyToSymbol
(a,host_a,Nbytes);
cudaMemcpyToSymbol
(b,host_b,Nbytes);
init<<< NB, NT>>>(c);
add<<< NB, NT>>>(c);
return 0;
}
vectoradd_constant.cu
先端GPGPUシミュレーション工学特論
2015/05/14
実行時間


入力配列サイズ N = 213
スレッド数 NT = 256

409
カーネル
実行時間 [ms]
vectoradd
7.65×10‐3
vectoradd_constant
1.01×10‐2
各スレッドがコンスタントメモリの異なるアドレスにアクセス
すると，グローバルメモリよりも遅くなる
先端GPGPUシミュレーション工学特論
2015/05/14
コンスタントメモリ（同一アドレス参照）
#define N (8*1024) //64kBに収める
#define Nbytes (N*sizeof(float))
#define NT (256)
#define NB (N/NT)
__constant__ float a, b;
__global__ void init(float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
c[i] = 0.0f;
}
__global__ void add(float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
c[i] = a + b;
}
int main(void){
float *c;
float host_a,host_b;
410
host_a=1.0f;
host_b=2.0f;
cudaMalloc((void **)&c,Nbytes);
//host_a,host_bが配列ではないので
//アドレスを取り出すために&を付ける
cudaMemcpyToSymbol
(a,&host_a,sizeof(float));
cudaMemcpyToSymbol
(b,&host_b,sizeof(float));
init<<< NB, NT>>>(c);
add<<< NB, NT>>>(c);
return 0;
}
vectoradd_broadcast.cu
先端GPGPUシミュレーション工学特論
2015/05/14
実行時間


入力配列サイズ N = 213
スレッド数 NT = 256

411
カーネル
実行時間 [ms]
vectoradd
7.65×10‐3
vectoradd_constant
1.01×10‐2
vectoradd_broadcast
7.55×10‐3
各スレッドがコンスタントメモリの同一アドレスにアクセスす
ると高速化
先端GPGPUシミュレーション工学特論
2015/05/14
実行時間


入力配列サイズ N = 220
スレッド数 NT = 256


412
カーネル
実行時間 [ms]
vectoradd
0.116
vectoradd_broadcast
0.042
配列サイズが多くなると，コンスタントメモリの利用が有効
定数を参照する場合，#defineで定義した方が高速に実
行できるので使いどころが難しい
 畳み込みのような処理を行うカーネル内で，全スレッドが
同じ係数テーブルにアクセスするような場合に有効
先端GPGPUシミュレーション工学特論
2015/05/14
ホストメモリへの効率的なアクセス



ホスト（CPU）とデバイス
（GPU）のやりとり
GPUでの処理を高速化し
続けると，ホスト‐デバイス
間のメモリコピーがボトル
ネック化
ホスト－デバイス間の転
送の高速化
GPU
Chip
SM
SM
共有
メモリ
L1キャッ
シュ
レジ
スタ
レジ
スタ
レジ
スタ
レジ
スタ
CUDA CUDA CUDA CUDA Core Core Core Core
L1キャッ
シュ
レジ
スタ
レジ
スタ
共有
メモリ
レジ
スタ
レジ
スタ
CUDA CUDA CUDA CUDA Core Core Core Core
L2キャッシュ
ローカルローカル
メモリ
メモリ
･･･
ローカルローカル
メモリ
メモリ
･･･
コンスタントメモリ
ホスト
メモリ
テクスチャメモリ
グローバルメモリ
413
先端GPGPUシミュレーション工学特論
2015/05/14
ページロック（ピン）メモリ

OSによって管理されるメモリのうち，ページアウトしな
いことが保証されているメモリ

OSが配列の物理アドレスを記憶

GPUが物理アドレスを知れば，ダイレクトメモリアクセ
ス（DMA）を使ってホストとデータを交換できる
414
先端GPGPUシミュレーション工学特論
2015/05/14
OSによる記憶管理

記憶管理


メモリのアドレスをプロセス固有の仮想アドレスに変換して
割り付け
仮想記憶方式を採用


415
メモリの物理的な大きさに依存せず，また不連続なメモリ領
域を連続に見せかける方式
1個の記憶装置を占有しているようにプログラム可能
先端GPGPUシミュレーション工学特論
2015/05/14
仮想記憶方式

仮想的な記憶装置上のアドレス


メモリ上のアドレス


仮想アドレス，論理アドレス
仮想アドレスと実アドレス
の変換はOSが管理
実アドレス，物理アドレス
多重仮想記憶



416
システム内に複数の仮想アドレス空間を形成し，プロセス
ごとに割り当て
現在の計算機の主流
CPUとGPUも仮想アドレス空間が異なる
先端GPGPUシミュレーション工学特論
2015/05/14
ページング方式

仮想メモリ空間と物理メモリ空間を一定サイズの
ページと呼ばれる単位に分割して管理

ページテーブル


仮想アドレスから実アドレスの対応表
仮想アドレスから実アドレスへの変換はページ単位で実行
CPU
_____
_____
_____
_____
_____
_____
_____
_____
_____
メモリ
_____
_____
_____
417
ページ
テーブル
_____
_____
_____
プロセスA
ページ
OS
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
先端GPGPUシミュレーション工学特論
スレッド
メモリ
2015/05/14
ページアウト

物理メモリ上にない仮想メモリを参照



補助記憶装置（ハードディスクなど）に退避されたデータ
ページフォルトという割り込みがかかり，OSに制御が移行
ページフォルトがおこると膨大な時間がかかる


OSは物理メモリ上のアドレスを追い出す（ページアウト）
必要なページを補助記憶装置から物理メモリ上に読み込む
ハード
ディスク
419
CPU
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
メモリ
_____
_____
_____
_____
_____
ページ
テーブル
_____
_____
_____
_____
プロセスA
ページ
OS
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
先端GPGPUシミュレーション工学特論
スレッド
メモリ
ページアウト
2015/05/14
ページロック（ピン）メモリ

ページアウトされない事がOSによって保証




ピンで刺された紙のように固定されている
pinned memory
そのメモリの物理アドレスにアクセスしても安全
全てのメモリをページロックにすると，他のプログラ
ムが起動できなくなる可能性がある
CPU
_____
_____
_____
_____
_____
_____
_____
_____
_____
メモリ
_____
_____
_____
420
ページ
テーブル
_____
_____
_____
プロセスA
ページ
OS
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
_____
先端GPGPUシミュレーション工学特論
スレッド
メモリ
2015/05/14
CUDAによるメモリ転送

ページング可能なメモリを使ってもDMA転送を行う
1.
2.

ページング可能なシステムバッファからページロックされ
たステージングバッファへコピー
ステージングバッファからGPUへDMAでコピー
ページロックメモリを使う事で約2倍の転送速度が
期待される

421
Fermi世代のGPUでは，10MBを超えるデータを転送すると
きにページロックメモリを使用すると効率的
先端GPGPUシミュレーション工学特論
2015/05/14
ページロックホストメモリの確保と解放

cudaMallocHost()




ページロックされたホストメモリを確保
cudaMallocHost((void **)&ホスト変数名, サイズ);
cudaMallocHostで確保されたメモリはmallocで確保さ
れたメモリと同様に利用可能
cudaFreeHost()


422
cudaMallocHostで確保されたホストメモリを解放
cudaFreeHost(ホスト変数のアドレス);
先端GPGPUシミュレーション工学特論
2015/05/14
ページロックホストメモリの確保
（もう一つの方法）

cudaHostAlloc()


cudaHostAlloc((void **)&ホスト変数名,サイズ,フラグ);
フラグ
 cudaHostAllocDefault



cudaHostAlloc()の働きをcudaMallocHost()と同じにする
確保したメモリの解放はcudaFreeHost()を利用
後述のゼロコピーホストメモリを確保するために利用


423
cudaHostAllocDefault以外のフラグを指定
 フラグについては後述
cudaHostAllocで統一した方が楽かもしれない
先端GPGPUシミュレーション工学特論
2015/05/14
ページング可能メモリを使ったコピー
#include<stdio.h>
#include<stdlib.h>
#define N (1024*1024*64)
#define Nbytes (N*sizeof(float))
cudaMemcpyHostToDevice);
//GPU→CPUはコメントを外す
//cudaMemcpy(data,dev_data,Nbytes, // cudaMemcpyDeviceToHost);
int main(){
float *data,*dev_data;
cudaEvent_t start,stop;
float elapsed_time_ms = 0.0f;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime
(&elapsed_time_ms, start,stop);
printf("%e ms¥n",elapsed_time_ms);
cudaEventDestroy(start);
cudaEventDestroy(stop);
free(data);
cudaFree(dev_data);
return 0;
data = (float *)malloc(Nbytes);
cudaMalloc( (void **)&dev_data, Nbytes);
}
cudaEventRecord(start, 0);
cudaMemcpy(dev_data,data,Nbytes, 424
copy_pagable.cu
先端GPGPUシミュレーション工学特論
2015/05/14
ページロックホストメモリを使ったコピー
#include<stdio.h>
#include<stdlib.h>
#define N (1024*1024*64)
#define Nbytes (N*sizeof(float))
cudaMemcpyHostToDevice);
//GPU→CPUはコメントを外す
//cudaMemcpy(data,dev_data,Nbytes, // cudaMemcpyDeviceToHost);
int main(){
float *data,*dev_data;
cudaEvent_t start,stop;
float elapsed_time_ms = 0.0f;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime
(&elapsed_time_ms, start,stop);
printf("%e ms¥n",elapsed_time_ms);
cudaEventDestroy(start);
cudaEventDestroy(stop);
cudaFreeHost(data);
cudaFree(dev_data);
return 0;
cudaHostAlloc( (void **)&data,
Nbytes, cudaHostAllocDefault);
cudaMalloc( (void **)&dev_data, Nbytes);
}
cudaEventRecord(start, 0);
cudaMemcpy(dev_data,data,Nbytes, 425
copy_pagelocked.cu
先端GPGPUシミュレーション工学特論
2015/05/14
データ転送の性能

入力配列サイズ N = 226
メモリ
実行時間 [ms]
転送速度 [GB/s]
CPU to GPU / GPU to CPU CPU to GPU / GPU to CPU
ページング可能
99.1 / 117
2.5 / 2.1
ページロック
44.8 / 42.1
5.6 / 5.9

ページロックメモリの利用により転送速度が約2倍に向上
相対的に遅かったGPU→CPUの転送が特に高速化
426
先端GPGPUシミュレーション工学特論

2015/05/14
CUDA4.0以降でのページロックメモリ

実用的にはCUDA4.1以降


CUDA4.0では適用できるメモリに制約がある
mallocで宣言してcudaHostRegisterでページロッ
クメモリとして割当


cudaHostRegister(ホスト変数アドレス, サイズ, フラグ)
フラグ





427
cudaHostRegisterDefault
cudaHostRegisterPortable
cudaHostRegisterMapped
ページロックにするだけならcudaHostRegisterDefault
cudaHostUnregisterによって割当を解除
先端GPGPUシミュレーション工学特論
2015/05/14
ページング可能メモリを使ったコピー
#include<stdio.h>
#include<stdlib.h>
#define N (1024*1024*64)
#define Nbytes (N*sizeof(float))
#define ByteToGByte (1.0/(1024*1024*1024))
#define SecToMillisec (1.0/1000)
cudaEventRecord(start, 0);
cudaMemcpy(dev_data,data,Nbytes, cudaMemcpyHostToDevice);
//cudaMemcpy(data,dev_data,Nbytes, // cudaMemcpyDeviceToHost);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime
(&elapsed_time_ms, start,stop);
printf("%e ms¥n",elapsed_time_ms);
int main(){
float *data,*dev_data;
cudaEvent_t start,stop;
float elapsed_time_ms = 0.0f;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaEventDestroy(start);
cudaEventDestroy(stop);
free(data);
cudaFree(dev_data);
return 0;
data = (float *)malloc(Nbytes);
cudaMalloc( (void **)&dev_data,
Nbytes);
}
copy_pagable.cu
428
先端GPGPUシミュレーション工学特論
2015/05/14
HostRegisterによる割当
#include<stdio.h>
#include<stdlib.h>
#define N (1024*1024*64)
#define Nbytes (N*sizeof(float))
#define ByteToGByte (1.0/(1024*1024*1024))
#define SecToMillisec (1.0/1000)
int main(){
float *data,*dev_data;
cudaEvent_t start,stop;
float elapsed_time_ms = 0.0f;
cudaEventCreate(&start);
cudaEventCreate(&stop);
data = (float *)malloc(Nbytes);
cudaMalloc( (void **)&dev_data,
Nbytes);
cudaHostRegister((void **)&data, }
Nbytes,cudaHostRegisterDefault);
429
cudaEventRecord(start, 0);
cudaMemcpy(dev_data,data,Nbytes, cudaMemcpyHostToDevice);
//cudaMemcpy(data,dev_data,Nbytes, // cudaMemcpyDeviceToHost);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaEventElapsedTime
(&elapsed_time_ms, start,stop);
printf("%e ms¥n",elapsed_time_ms);
cudaEventDestroy(start);
cudaEventDestroy(stop);
free(data);
cudaFree(dev_data);
return 0;
copy_hostregister.cu
先端GPGPUシミュレーション工学特論
2015/05/14
データ転送の性能

入力配列サイズ N = 226
メモリ
ページング可能
実行時間 [ms]
転送速度 [GB/s]
CPU to GPU / GPU to CPU CPU to GPU / GPU to CPU
99.1 / 117
2.5 / 2.1
cudaHostRegister
Default
76.4 / 85.2
3.3 / 2.9
cudaHostRegister
Portable
81.0 / 125 3.1 / 2.0
cudaHostRegister
Mapped
93.1 / 122
2.7 / 2.1

430
フラグによって性能が変わるが，どれも大して高速化しない
先端GPGPUシミュレーション工学特論
2015/05/14
ゼロコピーホストメモリ

ページロックホストメモリを使うとメモリがページアウ
トされず，アドレスが固定



GPUからOSを介さずCPUのメモリをコピーできた
GPUから書き込む事はできないのか？
デバイスからホストメモリを直接読み書きできるよう
アドレスをマッピングすれば可能


431
CPU‐GPU間のメモリコピーを要求しない
ゼロコピー
 GPUにおける根本的な制約「GPUはホストメモリを直接参
照できない」を回避
先端GPGPUシミュレーション工学特論
2015/05/14
ゼロコピーメモリ利用の可否

cudaGetDevicePropertiesを利用


使い方はcudaSetDeviceの時と同じ
cudaDeviceProp型構造体のメンバcanMapHostMemory
を参照
#include<stdio.h>
int main(void){
int deviceCount = 0,dev;
cudaDeviceProp deviceProp;
//ゼロコピーメモリが利用できるなら
//supportと表示
if(deviceProp.canMapHostMemory==1)
printf("supports¥n");
else
printf("doesn't support¥n");
//GPUの数を確認
cudaGetDeviceCount(&deviceCount);
for(dev=0;dev<deviceCount;dev++){
//情報を取得するGPUの選択
cudaSetDevice(dev);
cudaGetDeviceProperties
(&deviceProp, dev);
432
}
return 0;
}
先端GPGPUシミュレーション工学特論
2015/05/14
ゼロコピーメモリの確保

cudaHostAlloc()


cudaHostAlloc((void **)&ホスト変数名,サイズ,フラグ);
フラグ（|（or）で複数指定可能）
 cudaHostAllocMapped


GPUのアドレス空間にマッピングされたホストメモリを宣言
cudaHostAllocWriteCombined
ホストメモリのキャッシュ管理をOSから切り離す
 PCI‐Ex経由のデータ転送を高速化できる可能性がある
 ホストからの読込は非常に低速
 ホストが書き，GPUが読むバッファとしての利用に適する


433
cudaHostAllocPortable
先端GPGPUシミュレーション工学特論
2015/05/14
ホストメモリとデバイスメモリの対応付け

cudaHostAlloc()はCPUのポインタを返す

CPUとGPUでは仮想メモリ空間が異なる


CPUのポインタからGPUで利用できるポインタを取得する必
要がある
cudaHostGetDevicePointer()


cudaHostGetDevicePointer((void **)デバイス変数, (void *)ホストポインタ変数, 0);
最後の0はとりあえず入れておかなければならない

ドキュメント*にも"At present, Flags must be set to 0."と
書かれている
*http://docs.nvidia.com/cuda/samples/0_Simple/simpleZeroCopy/doc/CUDA2.2PinnedMemoryAPIs.pdf
434
先端GPGPUシミュレーション工学特論
2015/05/14
その他の準備

複数のGPUを搭載している場合は利用するGPUを指定


ホストメモリをデバイスから直接読み書きできるよう
アドレスのマッピングを許可


cudaSetDevice();
cudaSetDeviceFlags(cudaDeviceMapHost);
CUDAの関数を実行する前に上二つの命令を呼ぶ

435
呼ばないとcudaHostGetDevicePointer()がエラーを返す
先端GPGPUシミュレーション工学特論
2015/05/14
初期化と転送も含めたベクトル和
#define N (1024*1024)
#define Nbytes (N*sizeof(float))
#define NT (256)
#define NB (N/NT)
void init(float *a, float *b, float *c){
int i;
for(i=0;i<N;i++){
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
}
}
__global__ void add(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
float *dev_a,*dev_b,*dev_c;
cudaEvent_t start,stop;
float elapsed_time_ms = 0.0f;
cudaEventCreate(&start);
cudaEventCreate(&stop);
cudaHostAlloc((void **)&a, Nbytes,
cudaHostAllocDefault);
cudaHostAlloc((void **)&b, Nbytes, cudaHostAllocDefault);
cudaHostAlloc((void **)&c, Nbytes, cudaHostAllocDefault);
cudaMalloc((void **)&dev_a, Nbytes);
cudaMalloc((void **)&dev_b, Nbytes);
cudaMalloc((void **)&dev_c, Nbytes);
c[i] = a[i] + b[i];
}
int main(){
float *a,*b,*c;
436
vectoradd_host.cu
先端GPGPUシミュレーション工学特論
2015/05/14
初期化と転送も含めたベクトル和
printf("%e ms¥n",elapsed_time_ms);
cudaEventRecord(start, 0);
//ホストで初期化
init(a,b,c);
//転送
cudaMemcpy(dev_a, a, Nbytes, cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, Nbytes, cudaMemcpyHostToDevice);
cudaMemcpy(dev_c, c, Nbytes, cudaMemcpyHostToDevice); }
//デバイスでベクトル和
add<<<NB,NT>>>(dev_a,dev_b,dev_c);
//転送
cudaMemcpy(c, dev_c, Nbytes, cudaMemcpyDeviceToHost);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaDeviceSynchronize();
cudaEventElapsedTime
(&elapsed_time_ms, start,stop);
437
cudaFreeHost(a);
cudaFreeHost(b);
cudaFreeHost(c);
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
return 0;
vectoradd_host.cu
先端GPGPUシミュレーション工学特論
2015/05/14
ゼロコピーによるベクトル和
#define N (1024*1024)
#define Nbytes (N*sizeof(float))
#define NT (256)
#define NB (N/NT)
cudaEvent_t start,stop;
float elapsed_time_ms = 0.0f;
cudaEventCreate(&start);
cudaEventCreate(&stop);
__global__
void init(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
a[i] = 1.0;
b[i] = 2.0;
c[i] = 0.0;
}
__global__ void add(float *a, float *b, float *c){
int i = blockIdx.x*blockDim.x
+ threadIdx.x;
c[i] = a[i] + b[i];
}
cudaSetDevice(0);
cudaSetDeviceFlags(cudaDeviceMapHost);
int main(){
float *a,*b,*c;
float *dev_a,*dev_b,*dev_c;
438
cudaHostAlloc( (void **)&a, Nbytes, cudaHostAllocWriteCombined | cudaHostAllocMapped);
cudaHostAlloc( (void **)&b, Nbytes, cudaHostAllocWriteCombined |
cudaHostAllocMapped);
cudaHostAlloc( (void **)&c, Nbytes, cudaHostAllocWriteCombined |
cudaHostAllocMapped);
vectoradd_zerocopy.cu
先端GPGPUシミュレーション工学特論
2015/05/14
ゼロコピーによるベクトル和
//CPUとGPUのアドレスをマッピング
cudaHostGetDevicePointer(&dev_a,a,0);
cudaHostGetDevicePointer(&dev_b,b,0);
cudaHostGetDevicePointer(&dev_c,c,0);
}
cudaEventRecord(start, 0);
//転送無しでカーネルから読み書き
init<<<NB,NT>>>(dev_a,dev_b,dev_c);
add <<<NB,NT>>>(dev_a,dev_b,dev_c);
cudaEventRecord(stop, 0);
cudaEventSynchronize(stop);
cudaDeviceSynchronize();
cudaEventElapsedTime
(&elapsed_time_ms, start,stop);
printf("%e ms¥n",elapsed_time_ms);
cudaFreeHost(a);
cudaFreeHost(b);
cudaFreeHost(c);
return 0;
439
vectoradd_zerocopy.cu
先端GPGPUシミュレーション工学特論
2015/05/14
実行時間（初期化とベクトル和＋転送）


入力配列サイズ N = 220
スレッド数 NT = 256


440
カーネル
実行時間 [ms]
全てCPUで実行
8.74
全てGPUで実行
0.249
vectoradd_host
7.88
vectoradd_zerocopy
3.44
−ベクトル和
+転送
ゼロコピーは全てCPUで実行するよりは早い
全てGPUで実行するよりかなり遅い
先端GPGPUシミュレーション工学特論
2015/05/14
ゼロコピーの使いどころ

演算量の多いカーネル



通常はデータ転送→カーネル実行
ゼロコピーはカーネルの実行中にデータを転送
既にコピーされたデータで大量の計算を行うことでデータの
コピーにかかる時間を隠蔽

GPUのメモリ利用の節約

制約

GPUからホストメモリにアクセスしてもキャッシュされない
複数回の読み書きではPCI‐Ex経由の転送が複数発生
441
先端GPGPUシミュレーション工学特論

2015/05/14
ゼロコピーの使いどころ

ノートPC等でGPUがシステムのチップセットに組み込ま
れている場合*



GPUとCPUがメインメモリを物理的に共有
ゼロコピーを使うと常にパフォーマンスが向上
 ページロックメモリを利用しすぎるとシステムの性能が低下
GPUがチップセットに組み込まれているかの確認


cudaGetDeviceProperties()を利用
cudaDeviceProp型構造体のメンバintegratedを参照
 trueならintegrated GPU, falseならdiscrete GPU
*最近はこういう製品がない
442
先端GPGPUシミュレーション工学特論
2015/05/14

Download Report