約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のソース読んでみようかな。

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

__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 を使ったコードの主要部分

#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";
}