DEIM Forum 2023 2a-7-2

GPU 用データ並列プリミティブの大規模データへの拡張
中野 雅史†

三浦 真矢†

常

穹†

宮崎

純†

† 東京工業大学情報理工学院情報工学系 〒 152–8550 東京都目黒区大岡山二丁目 12 番 1 号
E-mail: †{nakano,miura}@lsc.c.titech.ac.jp, ††{q.chang,miyazaki}@c.titech.ac.jp
あらまし

本研究では，既存の GPU 用データ並列プリミティブを使用し，CPU と GPU を併用して GPU のメモリ

サイズを超えるデータを処理する手法を提案する．提案手法では，GPU 上の計算と CPU-GPU 間のデータ転送のパイプライン化とストリームの使用，および CPU 上の計算と GPU に関する処理のオーバーラップにより，処理全体の効率化を図る．評価実験では，比較手法として，CPU 上で POSIX スレッドを用いてデータ並列プリミティブと同様の処理を実装し，実行時間を比較した．また，プロファイラを使用して GPU 上での処理の様子を調べた．評価実験の結果から，本研究で扱った全てのデータ並列プリミティブにおいて，比較手法よりも高速な処理が可能であることが判明した．
キーワード 先進ハードウェア活用・GPU, 大規模データ処理, データ並列プリミティブ
法では，GPU 上での計算と CPU-GPU 間のデータ転送をパイ

1 はじめに社会で収集・集積されるデータは年々増加しており，それらは様々な用途に利用される．特に，地震や気象の予測・シミュレー
ションやディープラーニングなどの高度な技術には，膨大な計算量やデータ量が必要とされる場合が多い．大規模データに対する処理の高速化や効率化には，高性能なプロセッサの開発のようなハードウェア的手法のほか，既存のハードウェアを組み合わせて高性能なアーキテクチャを構成する手法がある．後者の手法の 1 つに，ドメインスペシフィックなプロセッサを他のドメインや汎用計算に活用する，GPGPU（General-purpose

computing on Graphics Processing Units）がある．

プライン化して，CUDA Stream を用いてオーバーラップさせる効率化と，CPU 上の計算と GPU に関する処理をオーバーラップさせる効率化を行う．評価実験では，比較手法として，

CPU 上で POSIX スレッドを用いてデータ並列プリミティブと同様の処理を実装し，実行時間を比較した．評価実験の結果か
ら，本研究で扱った reduce，hashtable，segmented sort の全てで，比較手法よりも高速な処理が可能であることが判明した．

2 GPGPU
画像処理を目的とする GPU は，単純な計算処理を行う数千個のコアを使った並列処理が特徴である．GPGPU は，この特GPGPU は，画像処理を目的として設計されたプロセッサ

徴を汎用計算に利用する技術であり，GPU の構造に適したプである GPU(Graphics Processing Unit) を，汎用計算に利用

ログラムを動作させた場合，汎用計算の逐次処理を得意とするする技術である．汎用的な処理に適した CPU に対し，単純CPU よりも高速な処理が期待できる．

な設計のコアを多数搭載する GPU は並列計算に優れており，

GPGPU では，GPU の性能を最大限に活かすプログラムを

GPGPU では，GPU 上での並列処理を組み合わせて高速な汎作成することが重要である．GPGPU プログラムの処理の基本

用計算を実現する．また，入力配列に対する scan [1] や sort [2]，的な流れを次に示す．

reduce などの汎用計算を GPU 上で実装したものとして，デー（ 1 ） GPU で扱うデータのために GPU 上のメモリ（以下，タ並列プリミティブがある [3]．

デバイスメモリと呼ぶ）を確保する

デ ー タ 並 列 プ リ ミ ティブ を 実 装 す る ラ イ ブ ラ リ と し て ，（ 2 ） CPU から GPU にデータを転送する

CUDPP [4] や ModernGPU [5] などが存在する．これらのライ（ 3 ） GPU で処理を行う関数（以下，カーネルと呼ぶ）をブラリは，NVIDIA 社の CUDA のような GPGPU のための起動し，処理を実行する開発環境で使用することができる．しかし，既存のデータ並列（ 4 ） GPU から CPU に計算結果のデータを転送する

プリミティブは，入力配列として GPU のメモリサイズを超え本研究では，NVIDIA 社の GPU と GPGPU 用の開発環境でないデータを対象としているため，処理できるデータの大きさある CUDA を用いて提案手法の実装と評価実験を行うため，に制限がある．既存の GPU の多くは，CPU のメモリよりも以下では CUDA を用いたプログラミングについて説明する．

小さなメモリを持つため，この制限は大規模データの処理を行う際に問題となる．
本研究では，データの分割と CPU 上での処理，および GPU

2. 1 CUDA
2. 1. 1 CUDA プログラムの処理の流れ上での既存のデータ並列プリミティブを使用した処理を組み合

CUDA プログラミングでは，GPU で処理するデータの入力

わせ，データ全体に対して演算を行う手法を提案する．提案手や生成，使用するデバイスメモリの確保と解放，CPU メモリ

（以下，ホストメモリと呼ぶ）とデバイスメモリの間のデータ転ものであり，提供しているライブラリには CUDPP や Mod-送などは，CPU 上で呼び出され実行される関数によって行わ

ernGPU などがある．本研究で使用したデータ並列プリミティ

れる．そのため，CUDA プログラムは，CPU 上で実行されるブについて以下に述べる．

ホストコードと，CUDA カーネルのように GPU 上で実行されるデバイスコードで構成される．NVIDIA の CUDA コンパイ
ラである nvcc（NVIDIA CUDA Compiler）は，コンパイル時にホストコードとデバイスコードを分離する．基本的な CUDA
プログラムの処理の流れは，2 で示した GPGPU プログラムの処理の流れと同様で，最後にデバイスメモリの解放を行うことが加わる．
（3）を除く４つの操作は，CUDA Runtime API が提供する関数で行うことができる．また，（3）の操作で実行されるカーネルは，デバイスコードに記述された CUDA カーネルである。

reduceこ の デ ー タ 並 列 プ リ ミ ティブ は ，要 素 数 n の 配 列[a0 , a1 , ..., an−1 ] と二項演算子 ⊕ を受け取り，a0 ⊕a1 ⊕...⊕an−1
を計算して出力する．

hashtable
このデータ並列プリミティブは，GPU 上でハッシュテーブル

[6] を実装したものである。CUDPP の hashtable は，入力として同じ長さの key と value の配列が与えられると，key 配列から重複する key を取り除いてソートし，配列 [key0 , key1 , ..., keyn−1 ]2. 1. 2 CUDA ストリーム
一般にデバイスメモリサイズは，ホストメモリサイズよりも小さい場合が多い．したがって，GPU を使用する処理では，
デバイスメモリより大きなデータを処理する場合，両者の間のを作成する．その後，この配列の各 key とそのインデックスを，式 (1) のハッシュ値を用いてハッシュテーブルに格納する．一方で，value 配列は key の順にソートして保持される．

データ転送を何度も行う必要がある．これは処理の実行時間に大きく影響し，ボトルネックとなる可能性がある．

ハッシュ値 =

そこで，CUDA は，マルチストリームによるデータの非同期((random 定数 1ˆkey) + random 定数 2)%定数素数処理を実現する API を提供している．CUDA Runtime API の(1)

cudaStreamCreate( ) 関数で生成したストリームを cudaMemcpyAsync( ) 関数や GPU カーネルの引数に渡すと，これらのある key に対して計算されたハッシュ値が，値の異なる他の key関数によるデータ転送と計算処理は，発行された順序に従い，に対して計算されたハッシュ値と一致する場合があり，この現

CUDA ランタイムが判断したタイミングで実行される。複数象を衝突と呼ぶ．CUDPP の hashtable は Cuckoo Hashing [7]のストリームを使用する場合，あるストリームに割り当てられを採用しており，衝突が起きたとき，既に格納されている方のたデータ転送とカーネルによる処理計算は，他のストリームに

key に対して random 定数 1 と random 定数 2 を異なる値に変割り当てられた異なる種類の処理とオーバーラップされる．まえてハッシュ値を再計算する，という操作を繰り返す．定めた

た，非同期なデータ転送を行う場合，ページングできない領域

試行回数を超えた場合，もう 1 つのハッシュテーブルへの格納（ピンメモリ）をホストメモリとして使用する必要があり，ピを試みる．

ンメモリの確保は cudaMallocHost( ) 関数を用いて行うことがsegmented sort

できる．ストリームを使用したパイプライン処理の様子を図 1に示す．

このデータ並列プリミティブは，配列 A と，m 個のセグメントの境界を定義する要素数 m − 1 の配列 B を受け取り，配列 A をセグメントごとにソートする．このとき，B[i], i = 0, 1, ..., m−2

Memcpy̲HtoD

Kernel Exce

Memcpy DtoH

Pipelining

は，i + 1 番目のセグメントの開始インデックスを表す．

merge sort
このデータ並列プリミティブは，配列と比較に用いる関数を

stream1
stream2

HtoD

受け取り、その関数に従う順に配列をソートする．

kernelDtoH

HtoD

kernelDtoH

HtoD

kernelstream3

3. 2 拡張データ並列プリミティブの実装
DtoH

三浦ら [8] は，既存のデータ並列プリミティブを用い，入力
データを分割して適切な順序で処理を行うことで，GPU のメモ

図 1 ストリームを用いたパイプライン処理の例

リサイズよりも大きな入力データの処理が可能な，transform，

reduce，scan，scatter，sort を実装した．CPU-GPU 間のデー
タ転送と計算をパイプライン化し，CUDA Stream を使用して

3 関連研究3. 1 データ並列プリミティブ
データ並列プリミティブは，汎用計算を GPU 上で実装した

それらをオーバーラップさせることで処理の効率化も行なった．
評価実験では，比較手法として CPU 上で各データ並列プリミ
ティブと同じ処理を実装し，scatter 以外のプリミティブで比較手法より高速な処理が可能であることを示した．

3. 3 データ並列プリミティブを活用したデータベース処理の実装

Input:

大原ら [9] は，データベースに蓄積されたシーケンスデータ

m

n

…

…

に対する行パターンマッチングを標準化した SQL/RPR（Row

Pattern Recognition）の処理に，データ並列プリミティブを使…

用して，GPU 上で処理する手法を提案した．SQL/RPR の主…

な処理は以下である．
（ 1 ） PARTITION BY 句を使用したデータベースのパー
ティション

GPU

CPU

multistream pipeline processing

multithread processing（ 2 ） ORDER BY 句を使用したデータの並び替え（ 3 ） DEFINE 句を使用した行パターン変数の暫定マッピ
ング（ 4 ） PATTERN 句を使用したパターンの抽出提案手法では（1）（，2）（，4）の処理の一部に，それぞれ CUDPPの hashtable，ModernGPU の segmented sort，ModernGPUの reduce を用いた．評価実験では，比較手法として PostgreSQL

Output:上で SQL/RPR を実装し，提案手法がクエリ処理の実行時間を短縮することを示した．

4 提案手法図 2 reduce の動作例

される値は uint2 型で，上位 32bit に key の値，下位 32bit に

本節では，既存のデータ並列プリミティブを使用して，GPUvalue の値を持つ．一方で，MultivalueHashTable は，与えらのメモリサイズよりも大きなデータを GPU と CPU を併用れた value 配列を key の順にソートして保持する．その取得して処理するための具体的な実装方法について説明する．本には cudppMultivalueHashGetAllValues() 関数を用い，cud-研究で扱ったデータ並列プリミティブは reduce，hashtable，ppHashRetrieve() 関数を使って key ごとのオフセットを取得

segmented sort である．

することができる．また，cudppHashRetrieve() 関数は，keyに対して求めたハッシュ値でテーブルを探索する．探索して見4. 1 reduce
はじめに，入力配列を m : n に分割し，それぞれ GPU 上と CPU 上で同時に処理を行う．GPU 上で処理する部分に対つけたエントリに格納されている key とクエリとして用いた

key を照合し，一致する場合には，その value を取得してオフセット情報の取得に利用する．

しては，配列を GPU のメモリサイズ以下に分割し，2.1.2 で各ストリームは，どちらも変更を加えた cudppMultivalue-

示した方法でパイプライン処理を行う．カーネルには CUDPP

HashGetAllValues() 関数と cudppHashRetrieve() 関数を呼びの reduce を使用する．CPU 上で処理する部分に対しては，出す．後者の関数によって得られる key ごとのオフセット配列

POSIX スレッドを用いて，配列をスレッド数と同数に分割してを用いて，前者の関数によって得られる key の順にソートされスレッドに割り当て，並列処理を行う．最後に GPU と CPU のた value 配列を，C++の std::unordered map を用いて実装し計算結果に対して CPU 上で計算を行い，出力する．また，事たハッシュテーブルに格納する．ハッシュテーブルへの格納は，前に行った実験でこれらの GPU 上と CPU 上での処理の速度

POSIX スレッドに key を割り当て，各スレッドがその key にを測定し，分割比 m，n をそれぞれ 0.65，0.35 とした．reduce

対応する部分を value 配列から取り出し，ハッシュテーブルへの動作例を図 2 に示す．

格納する方法で行う．この CPU 上での処理は，GPU の処理と

4. 2 hashtable
本研究では，GPU のメモリサイズより大きなデータを扱え

並行して行う．hashtbable の動作例を図 3 に示す．

4. 3 segmented sort

るように拡張した，データ並列プリミティブの hashtable を用

セグメントの大きさに従って処理方法を変えることで，効いてハッシュテーブル作成の高速化を図り，入力データを同一率的に入力配列内の全てのセグメントをソートする．GPU の

key ごとにグループ化する処理の効率化を目的とする．

使用可能なメモリサイズを M ，1 つの CUDA ストリームに

reduce における GPU 上での処理と同様に，key と value の割り当てたメモリサイズを M ′ とする．大きさが M ′ 以下の

2 つの入力配列を分割してストリームに割り当て，カーネルにセグメントは，それらを 1 つ以上集めてストリームに割り当

は CUDPP の MultivalueHashTable を用いる。

て，ModernGPU の segmented sort を適用する．M ′ より大MultivalueHashTable は，与えられた key 配列から重複す

きく M 以下のセグメントは，ストリームを使用したオーバー

る key を除いたユニークな key 配列を生成し，その配列のイ

ラップは行わず，1 つずつ順にデータ転送し ModernGPU の

ンデックスを value としてハッシュテーブルを作成する．格納Input:

key:

…

…

…

value:…

…

…

Input:

key:
value:

MemCpyHtoD

segmented sort
Partitioning

Kernel Execution

MemCpyHtoD

MemCpyDtoH

Kernel Execution

MemCpyHtoD

Make hashtable

MemCpyDtoH

Kernel Execution

Add to hashtable

MemCpyDtoH

hashtable

hashtable

segmented sort

GPU:
CPU:

Partitioning
merge sort

Add to hashtable

key:

GPU:

Output:Output:CPU:merge sorthashtable

図 4 segmented sort の動作例

図 3 hashtable の動作例

mergesort を適用する．M よりも大きなセグメントは，CPU
上でクイックソートの分割アルゴリズムにより M 以下に分割
し，ModernGPU の mergesort を適用する．
この分割では，ストリームを使用したオーバーラップによって短縮される時間よりも，分割にかかる時間の方が大きいため，

M ′ 以下への分割は行わない．分割アルゴリズム中の分割の基準となるピボットの選択は，分割する配列の長さを n としてインデックスが 0， n2 − 1，n − 1 の要素の中から中央値を求めてピボットとする方法を用いる．分割が終了しないことがないよ
うに，ひとつ前の分割で使用したピボットと同じ値がピボットに選択された場合，ピボット選択に使用する要素のインデック
スを 1， n2 ，n − 2 のようにずらせていく．また，CPU 上での分割は GPU 上の処理と並行して行うことで，全体の処理時間を短縮する．

GPU 上でソートしたセグメントをホストメモリにある入力
配列の適切な位置にコピーすることで，全てのストリームの処理が完了すると，入力配列はセグメントごとに key についてソートされたことになる．segmented sort の動作例を図 4 に示す．

value:CPU 上で実装した reduce は，入力配列の先頭から 2 つず
つ要素を取り出して計算を行い，全ての要素について計算し長さが半分になった配列に対し，同様の処理を再帰的に行う．

POSIX スレッドを使った実装では，入力配列をスレッド数で分割してスレッドに割り当て，最後に全てのスレッドの計算結果に対して，シングルスレッドで reduce を行う。POSIX スレッドを使った比較手法の実装を Algorithm1 に示す．

CPU 上で実装した hashtable は、std::unordered map をハッシュテーブルとして用い，格納される value は std::vector を用いて実装した．std::unordered map は木構造で実装される

std::map とは異なりハッシュ表で実装されており，ハッシュテーブルへの格納は，key と value の同じ長さの入力配列に対して同じインデックスの値をペアとして，先頭から 1 ペアずつ
ハッシュテーブルに追加する。POSIX スレッドを使った実装
では，スレッド数を N として以下の手順で処理を行う．
（ 1 ） 入力を M 個に分割する（ 2 ） i 番目の分割をさらに N 個に分割し N 個のハッシュテーブルを作成する（ 3 ） ハッシュテーブルを 2 つずつ結合する（ 4 ） ハッシュテーブルが 2 つになるまで（3）の処理を繰

5 評価実験
本節では，データ並列プリミティブと同様の処理を CPU 上り返す（ 5 ） i = M であれば，M × 2 個のハッシュテーブルから
ベースとするものを 1 つ選び，残りの M × 2 − 1 個を順にベーで実装した手法，および POSIX スレッドを用いて CPU 上でスのハッシュテーブルへ結合する

実装した手法を比較手法として，提案手法と比較し評価を行う．

（ 6 ） i = i + 1 として（2）に戻る

5. 1 比 較 手 法
比較手法として用いる，CPU 上でデータ並列プリミティブと同様の処理を実装した方法，および POSIX スレッドを用いて CPU 上で実装した方法について述べる．比較手法の実装には C++を用い，POSIX スレッドを使った実装ではスレッド数を 16 とした．

手順（1）と（5）は，実行時間への影響が少ない範囲でメモリ使用量を抑えるために行う．POSIX スレッドを使った比較手法の実装を Algorithm2 に示す．

CPU 上で実装した segmented sort は，入力配列のセグメントに対して順に std::sort を適用し，全てのセグメントをソートする．POSIX スレッドを使った実装では，はじめに要素数が n 以上のセグメントを C++の Boost ライブラリのAlgorithm 1 比較手法の reduce の擬似コード

Algorithm 2 比較手法の hashtable の擬似コード

1: function th reduce

1: function th maketable2:
3:

while num >= 1 dofor i ← 0 to num
do2
input[i] ← input[2 ∗ i] + input[2 ∗ i + 1]

4:

2:
3:
4:

for i ← 0 to num do
insert pair(key[i], values[i]) to tableend for5:

end for

5: end function

6:

num ← num
2

6:

7:

end while

7: function th mergetable

8:

result[tid] ← input[0]8:9: end function

9:

10:

10:

11: function reducefor i ← 0 to numKeys domove table src[i] to table dest[i]end for

11: end function

12:

inputSizenum ← numT
hreads12:

13:

for i ← 0 to numT hreads do

13: function hashtable

14:\\multi thread

14:

TH REDUCE()

15:

inputSizenum ← partition∗numT
hreadsfor j ← 0 to numT hreads do

16:

end for16:

17:\\single thread

17:

18:

TH REDUCE()

18:

19: end function

for i ← 0 to partition do

15:\\multi thread
TH MAKETABLE()

19:

end for20:

while numT able >= 2 do

21:
22:

able
for i ← 0 to numT
do2\\multi thread

block indirect sort を用いてソートする．次に行う要素数が n23:以下のセグメントのソートは，1 つのスレッドに 1 つのセグ

24:

end for

メントを割り当て，std::sort を適用する方法でソートを行う．

25:

able
numT able ← numT2

boost::block indirect sort の性能を調べる実験を事前に行い，26:

本研究では n を 1 × 10 とした．boost::block indirect sort は27:

end for

最悪計算時間と平均計算時間が n log n のソートアルゴリズムで

28:

for i ← 0 to pertition ∗ 2 − 1 do

7ある．POSIX スレッドを使った比較手法の実装を Algorithm3
に示す．

29:
30:
31:

T H M ERGET ABLE()

end while\\multi thread
T H M ERGET ABLE()
end for32: end function

5. 2 実 験 準 備
評価実験に用いたマシンの構成を表 1 に示す．
表1

実験に用いたマシンの構成を求める計算を行なった．hashtable は，一様分布に従う [0,CPUIntel Core i7-10700 (2.90GHz，8 コア)

9] の範囲の乱数で初期化した key 配列，要素数を n として [0,

RAM

24GB

OS

Ubuntu 20.04.5 LTS

n − 1] の範囲の一意の整数で初期化した value 配列を入力とNVIDIA GTX 1650 (Turing)
GPU

して受け取り，std::unordered map で実装した CPU 上のハッシュテーブルに全ての key と value のペアを格納するタスクを

CUDA コア

896

ベースクロック

1485MHz行なった．segmented sort は，一様分布に従う [1, 20] の範囲

メモリ量

4GB GDDR5の乱数で初期化した key 配列，hashtable と同様の value 配列，

CUDA

CUDA 11.7

および zipf 分布に従う [1, n − 1] の乱数で初期化した要素数

99 のオフセット配列を入力として受け取り，セグメントごと
5. 3 実 験 内 容
本実験では，比較手法と提案手法の実行時間の計測，およびNVIDIA Nsight Systems を用いて提案手法のプロファイルを行った．提案手法の実行時間の計測は，ホストメモリからデバ
イスメモリへのデータ転送開始前から，デバイスメモリからホ
ストメモリへの計算結果の転送と，CPU 上での必要な処理の両方が完了するまでを実行時間として計測した．

reduce は，二項演算子に + 用いて，配列の全ての要素の和に key 配列と value 配列を key についてソートするタスクを行なった．zipf 分布のパラメータ s は，1.5 とした．

reduce の入力には 32bit 浮動小数点型配列を用い，hashtableと segmented sort の key，value の入力配列には 32bit 整数型
配列を用いた．reduce では入力配列の長さを 0 から 3 × 109 (約11.2GB) まで変化させ，hashtable と segmented sort では入力
配列の長さを 0 から 1.6 × 109 (約 5.96GB) まで変化させた．

でも，カーネル実行が最も時間がかかる処理で，hashtable とAlgorithm 3 比較手法の segmented sort の擬似コード

同様な処理のオーバーラップが行われたことがわかる．

1: function th sort2:

sort(array)

7000

CPU

3: end function

5: function th multisort

multisort(array)

7: end function
8:
9: function segmented sort11:

for i ← 0 to numSegments do
if segmentSize >= n then

12:

TH MULTISORT()

10:

13:

end for15:

for i ← 0 to numSegments do

16:

if segmentSize < n then

17:
18:
19:

4000

3000

2000

1000

0

end if

14:

GPU

5000
elapsed time [ms]6:

CPUThread

6000

4:

0.0

0.2

0.4
0.6
number of elements

0.8

1.0
3e9

0.8

1.0
1.6e9

0.8

1.0
1.6e9

図 5 reduce の実行結果

TH SORT()end if
140000end forCPU

20: end function

CPUThread

120000GPU

5. 4 実 験 結 果
比較手法と提案手法の実行時間の測定結果を，図 5 から図 7

elapsed time [ms]100000に示す。reduce では，提案手法が，比較手法のうち CPU 実装に

80000

60000

40000

比べて約 6.09-6.84 倍，POSIX スレッドを使った CPU 実装に20000

比べて約 1.80-1.96 倍の性能，hashtable では，比較手法のうち

CPU 実装に比べて約 21.0-26.3 倍，POSIX スレッドを使った

0
0.0

0.2

CPU 実装に比べて約 2.76-3.84 倍の性能となった．segmented
sort では，比較手法のうち CPU 実装に比べて約 46.7-160.8 倍，

0.4
0.6
number of elements

図 6 hashtable の実行結果

POSIX スレッドを使った CPU 実装に比べて約 8.62-32.8 倍の性能となった．
提案手法のプロファイル結果を，図 8 から図 10 に示す。図1000000CPU
900000

ムに割り当てたメモリサイズ以下のセグメントに対して，GPU

800000上で行なった処理のプロファイル結果である．実験では 3 つま

700000

たは 4 つの CUDA ストリームを使い，図 8 から図 10 では，上部の黒枠で囲んだ部分は全てのストリームを合わせた処理の流れ，その下は各ストリームの処理の流れを表示している．図のelapsed time [ms]10 は，segmented sort の提案手法において，CUDA ストリー

500000400000300000青の部分はカーネル実行，赤の部分はデバイスメモリからホス

200000

トメモリへのデータ転送の処理を表している．reduce では，ホ100000

ストメモリからデバイスメモリへのデータコピーが実行時間の0

ン化とオーバーラップにより，カーネル実行とデバイスメモリ
からホストメモリへのデータコピーが隠蔽されたことがわかる．

hashtable では，カーネル実行が最も時間がかかる処理で，処理全体ではホストメモリからデバイスメモリへのデータコピー
が他の 2 つの処理に隠蔽されたことがわかる．segmented sortGPU

600000緑の部分はホストメモリからデバイスメモリへのデータ転送，

大半を占め，CUDA ストリームを利用した処理のパイプライ

CPUThread

0.0

0.2

0.4
0.6
number of elements

図 7 segmented sort の実行結果

同様の処理を CPU 上で実装した手法と，POSIX スレッドを用いて CPU 上で実装した手法を用意した．評価実験は，各デー
タ並列プリミティブに対して，入力配列の大きさを変えながら実行時間を測定した．評価実験の結果から，CPU 上の実装と
図 8 reduce のプロファイル結果

比べて約 6.09-160.8 倍，POSIX スレッドを使った CPU 上での実装と比べて約 1.80-32.8 倍の性能向上を実現した．
今後の課題として，提案手法に最適な実装方法を検討し，さらなる高速化を実現すること，CPU と GPU をより効果的に併用する手法の考案，本研究で対象としなかったデータ並列
プリミティブの拡張などが挙げられる．また，本研究で対象としたデータ並列プリミティブを用いて，大原ら [9] が提案した，

図 9 hashtable のプロファイル結果

GPU 上での SQL/RPR の処理手法のような既存のアプリケー
ションや処理手法を，メモリサイズの制約がないものに拡張することも 1 つの課題と考えられる．さらに，併用するプロセッ
サや計算機の性能差や通信手段からデータの分割比率を求める
自動最適化手法の考案も，本研究に関連する課題として挙げられる．

図 10 segmented sort のプロファイル結果

謝5. 5 考辞察

reduce では，提案手法において，入力配列を分割して CPUと GPU で同時に処理を行なったことに加えて，GPU 上での本研究は、JST CREST JPMJCR22M2 の支援を受けたものである。

処理のオーバーラップによってカーネル実行がデータ転送に隠蔽されたことで，実行時間が短縮されたと考えられる．その結果，提案手法は POSIX スレッドを使用して効率化を行なった
比較手法よりも実行時間が短くなったと考えられる．

hashtable では，提案手法において，ホストメモリからデバ
イスメモリへのデータ転送が，カーネル実行と逆方向のデータ転送に隠蔽されたことで，GPU 上での処理が効率的に行われた．さらに，GPU 上での計算結果を使用してハッシュテーブ
ルに要素を追加する CPU 上での処理を，GPU に関する処理とオーバーラップさせることで，POSIX スレッドを使用して効率化を行なった比較手法よりも高速に，ハッシュテーブルを作成することができたと考えられる．

segmented sort では，提案手法において，CUDA ストリー
ムに割り当てたメモリサイズ以下のセグメントに対して GPU
上で処理を行う部分で，hashtable と同様の効率化が行われた．
また，GPU のメモリサイズより大きなセグメントの数が比較的少ない場合には，CPU 上で行う分割処理の大部分が，GPU上で行うその他のセグメントの処理にオーバーラップする．これらにより，POSIX スレッドを使用して効率化を行なった比較手法に比べて，提案手法の性能が上回ったと考えられる．

6 終わりに本 研 究 で は ，デ ー タ 並 列 プ リ ミ ティブ で あ る reduce，

hashtable，segmented sort を，CPU と GPU を併用して GPUのメモリサイズを超えるデータの処理が可能となるように拡張する手法を提案した．
評価実験では，比較手法として，データ並列プリミティブと文

献[1] Shubhabrata Sengupta, Mark Harris, Michael Garland, andJohn D. Owens. Eﬃcient Parallel Scan Algorithms for
Many-core GPUs. eScholarship, University of California,
2011.
[2] Nadathur Satish, Mark Harris, and Michael Garland. Designing eﬃcient sorting algorithms for manycore GPUs. In
2009 IEEE International Symposium on Parallel & Distributed Processing, pp. 1–10, 2009[3] Bingsheng He, Ke Yang, Rui Fang, Mian Lu, Naga Govindaraju, Qiong Luo, and Pedro Sander. Relational joins ongraphics processors. In Proceedings of the 2008 ACM SIGMOD international conference on Management of data, pp.
511–524, June 2008.
[4] CUDPP, https://github.com/cudpp/cudpp (2023 年 1 月 9日アクセス)[5] Sean Baxter: moderngpu 2.0,
https://github.com/moderngpu/moderngpu (2023 年 1 月 9日アクセス)[6] Dan A Alcantara, Vasily Volkov, Shubhabrata Sengupta,
Michael Mitzenmacher, John D. Owens, and Nina Almenta.
Building an Eﬃcient Hash Table on the GPU. In GPU Computing Gems Jade Edition, pp. 39–53. Elsevier, 2012[7] R. Pagh and F. F. Rodler. Cuckoo hashing. Journal of Algorithms, vol. 51, no. 2, pp. 122–144, 2004.
[8] 三浦 真矢, 常 穹, 宮崎 純. GPU メモリサイズより大きなデー
タを対象としたデータ並列プリミティブの開発. 第 13 回デー
タ工学と情報マネジメントに関するフォーラム (DEIM2022),
B24-2(day2 p43), 2022.
[9] Tsubasa Ohara, Qiong Chang, and Jun Miyazaki. Fast
SQL/Row Pattern Recognition Query Processing Using
Parallel Primitives on GPUs. In Proceedings of the 32nd
International Conference on Database and Expert Systems Applications (DEXA2021), LNCS 12923, pp. 22-34,
September. 2021.