日曜研究室 LLVMによるプログラミング言語の実装 xv6ソースコードリーディング 技術 日記
日曜研究室

日曜研究室

まだ庭の小石の数を数え終わっていない

技術

std::sortと独自バイトニックソートとthrust::sortの速度

投稿日
2013/6/28 金曜日
カテゴリー
技術
タグ
, , ,
著者
Take it easy!

約100万要素（1024 * 1024)の32bit整数のソート速度を以下のプログラムで測ってみた。
全て C++で書き、64bit、O3でコンパイルした。
実行したマシンはCore 2 Duo 2GHz, Geforce 9400M。

  1. std::sortでソートするプログラム。GPU非使用。。コンパイラ: clang++。
  2. 前回Pythonで書いたバイトニックソートをCUDA(Driver API)にそのまま移植したもの。ホストコードコンパイラ: clang++、カーネルコンパイラ: nvcc。
  3. thrust::sortでソートするプログラム。CUDA Runtime API使用。コンパイラ: nvcc。

（ソースは本文の最後に載せるが、全部のせると長くなるので、バイトニックソートのカーネルと、thrust::sortを使ったプログラムの主要な部分のみのせる。）

結果は以下のとおり。

  1. std::sort 約0.37秒
  2. おれおれバイトニックソート 約1.08秒（うち0.01秒はCPU-GPU間のメモリコピー）
  3. thrust::sort 約0.05秒（うち0.01秒はCPU-GPU間のメモリコピー）

さすがにC++のstd::sortは、Pythonのsortedよりかなり速い。
おれおれバイトニックソートはGPU向けの最適化を一切してないので遅いのは仕方ない部分もあるんだけど、CPUのみで実行するstd::sortより遅いのはちょっとびっくり(*_*;
GPUが古くてメモリアクセスに対するコアレッシングが効きづらいからかなぁ。
で、thrust::sort超速い。
しかしthrust by CUDAはnvccじゃないとまともにコンパイルできないのはちょっとめんどくさい。
仕組み上仕方ないんだけど。

不思議なのは、おれおれバイトニックソートの方はO0〜O3で実行速度にほとんど変化がないのに対し、thrust::sortの方はO0とO3で30倍ぐらい速度が違う（O0だと1.8秒ぐらいかかる）こと。
気が向いたらthrustのソース読んでみようかな。

以下ソース。
おれおれバイトニックソートのカーネル

01
02
03
04
05
06
07
08
09
10
11
12
13
14
15
16
17
18
19
20
21
22
23
__global__ void bitonic_sort(int *data, int data_size, int chunk_size, int sub_chunk_size)
{
    int thread_index = blockDim.x * blockIdx.x + threadIdx.x;
 
    if (thread_index < data_size / 2) {
        int half_chunk_size = chunk_size / 2;
        int chunk_index = thread_index / half_chunk_size;
         
        int half_sub_chunk_size = sub_chunk_size / 2;
        int sub_chunk_index = thread_index / half_sub_chunk_size;
         
        bool up = (chunk_index % 2 == 0);
        int a = sub_chunk_size * sub_chunk_index + thread_index % half_sub_chunk_size;
        int b = a + half_sub_chunk_size;
         
        int va = data[a];
        int vb = data[b];
        if (va > vb == up) {
            data[a] = vb;
            data[b] = va;
        }
    }
}

thrust::sort を使ったコードの主要部分

01
02
03
04
05
06
07
08
09
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
#include <iostream>
#include <cstdlib>
#include "Timer.hpp"
#include <thrust/host_vector.h>
#include <thrust/device_vector.h>
#include <thrust/copy.h>
#include <thrust/sort.h>
 
void thrust_sort(int n, int randomseed)
{
    std::cout << "Start thrust_sort\n";
     
    // 時間計測タイマー
    Timer include_memcpy_timer, kernel_timer;
     
    // 元データの作成
    thrust::host_vector<int> hData;
    hData.reserve(n);
    srand(randomseed);
    for (int i = 0; i < n; i++) {
        double rnd = (double)rand() / RAND_MAX;
        hData.push_back(rnd * n);
    }
     
    // ホストからデバイスへコピー
    cudaDeviceSynchronize();
    include_memcpy_timer.Start();
    thrust::host_vector<int> dData(n);
    thrust::copy(hData.begin(), hData.end(), dData.begin());
     
    // ソート
    cudaDeviceSynchronize();
    kernel_timer.Start();
    thrust::sort(dData.begin(), dData.end());
    cudaDeviceSynchronize();
    kernel_timer.End();
     
    // デバイスからホストへコピー
    thrust::copy(dData.begin(), dData.end(), hData.begin());
    cudaDeviceSynchronize();
    include_memcpy_timer.End();
     
    // 結果の表示
    std::cout << "Time: " << kernel_timer.GetSeconds() << " sec. (include memcpy: " << include_memcpy_timer.GetSeconds() << " sec.)\n";
    std::cout << "Result: " << n << " elements\n";
    for (int i = 0; i < ((n < 128)?n:128); i++) {
        std::cout << hData[i] << ", ";
    }
    std::cout << "\nDone.\n";
}

コメントを残す

メールアドレスが公開されることはありません。 ※ が付いている欄は必須項目です



※画像をクリックして別の画像を表示

このサイトはスパムを低減するために Akismet を使っています。コメントデータの処理方法の詳細はこちらをご覧ください。