{"id":2678,"date":"2013-06-28T11:23:03","date_gmt":"2013-06-28T02:23:03","guid":{"rendered":"http:\/\/peta.okechan.net\/blog\/?p=2678"},"modified":"2013-06-28T11:23:03","modified_gmt":"2013-06-28T02:23:03","slug":"stdsort%e3%81%a8%e7%8b%ac%e8%87%aa%e3%83%90%e3%82%a4%e3%83%88%e3%83%8b%e3%83%83%e3%82%af%e3%82%bd%e3%83%bc%e3%83%88%e3%81%a8thrustsort%e3%81%ae%e9%80%9f%e5%ba%a6","status":"publish","type":"post","link":"https:\/\/peta.okechan.net\/blog\/archives\/2678","title":{"rendered":"std::sort\u3068\u72ec\u81ea\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u3068thrust::sort\u306e\u901f\u5ea6"},"content":{"rendered":"

\u7d04100\u4e07\u8981\u7d20\uff081024 * 1024)\u306e32bit\u6574\u6570\u306e\u30bd\u30fc\u30c8\u901f\u5ea6\u3092\u4ee5\u4e0b\u306e\u30d7\u30ed\u30b0\u30e9\u30e0\u3067\u6e2c\u3063\u3066\u307f\u305f\u3002
\n\u5168\u3066 C++\u3067\u66f8\u304d\u300164bit\u3001O3\u3067\u30b3\u30f3\u30d1\u30a4\u30eb\u3057\u305f\u3002
\n\u5b9f\u884c\u3057\u305f\u30de\u30b7\u30f3\u306fCore 2 Duo 2GHz, Geforce 9400M\u3002<\/p>\n

    \n
  1. std::sort\u3067\u30bd\u30fc\u30c8\u3059\u308b\u30d7\u30ed\u30b0\u30e9\u30e0\u3002GPU\u975e\u4f7f\u7528\u3002\u3002\u30b3\u30f3\u30d1\u30a4\u30e9: clang++\u3002<\/li>\n
  2. \u524d\u56dePython\u3067\u66f8\u3044\u305f\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8<\/a>\u3092CUDA(Driver API)\u306b\u305d\u306e\u307e\u307e\u79fb\u690d\u3057\u305f\u3082\u306e\u3002\u30db\u30b9\u30c8\u30b3\u30fc\u30c9\u30b3\u30f3\u30d1\u30a4\u30e9: clang++\u3001\u30ab\u30fc\u30cd\u30eb\u30b3\u30f3\u30d1\u30a4\u30e9: nvcc\u3002<\/li>\n
  3. thrust::sort\u3067\u30bd\u30fc\u30c8\u3059\u308b\u30d7\u30ed\u30b0\u30e9\u30e0\u3002CUDA Runtime API\u4f7f\u7528\u3002\u30b3\u30f3\u30d1\u30a4\u30e9: nvcc\u3002<\/li>\n<\/ol>\n

    \uff08\u30bd\u30fc\u30b9\u306f\u672c\u6587\u306e\u6700\u5f8c\u306b\u8f09\u305b\u308b\u304c\u3001\u5168\u90e8\u306e\u305b\u308b\u3068\u9577\u304f\u306a\u308b\u306e\u3067\u3001\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u306e\u30ab\u30fc\u30cd\u30eb\u3068\u3001thrust::sort\u3092\u4f7f\u3063\u305f\u30d7\u30ed\u30b0\u30e9\u30e0\u306e\u4e3b\u8981\u306a\u90e8\u5206\u306e\u307f\u306e\u305b\u308b\u3002\uff09<\/p>\n

    \u7d50\u679c\u306f\u4ee5\u4e0b\u306e\u3068\u304a\u308a\u3002<\/p>\n

      \n
    1. std::sort \u7d040.37\u79d2<\/li>\n
    2. \u304a\u308c\u304a\u308c\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8 \u7d041.08\u79d2\uff08\u3046\u30610.01\u79d2\u306fCPU-GPU\u9593\u306e\u30e1\u30e2\u30ea\u30b3\u30d4\u30fc\uff09<\/li>\n
    3. thrust::sort \u7d040.05\u79d2\uff08\u3046\u30610.01\u79d2\u306fCPU-GPU\u9593\u306e\u30e1\u30e2\u30ea\u30b3\u30d4\u30fc\uff09<\/li>\n<\/ol>\n

      \u3055\u3059\u304c\u306bC++\u306estd::sort\u306f\u3001Python\u306esorted\u3088\u308a\u304b\u306a\u308a\u901f\u3044\u3002
      \n\u304a\u308c\u304a\u308c\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u306fGPU\u5411\u3051\u306e\u6700\u9069\u5316\u3092\u4e00\u5207\u3057\u3066\u306a\u3044\u306e\u3067\u9045\u3044\u306e\u306f\u4ed5\u65b9\u306a\u3044\u90e8\u5206\u3082\u3042\u308b\u3093\u3060\u3051\u3069\u3001CPU\u306e\u307f\u3067\u5b9f\u884c\u3059\u308bstd::sort\u3088\u308a\u9045\u3044\u306e\u306f\u3061\u3087\u3063\u3068\u3073\u3063\u304f\u308a(*_*;
      \nGPU\u304c\u53e4\u304f\u3066\u30e1\u30e2\u30ea\u30a2\u30af\u30bb\u30b9\u306b\u5bfe\u3059\u308b\u30b3\u30a2\u30ec\u30c3\u30b7\u30f3\u30b0\u304c\u52b9\u304d\u3065\u3089\u3044\u304b\u3089\u304b\u306a\u3041\u3002
      \n\u3067\u3001thrust::sort\u8d85\u901f\u3044\u3002
      \n\u3057\u304b\u3057thrust by CUDA\u306fnvcc\u3058\u3083\u306a\u3044\u3068\u307e\u3068\u3082\u306b\u30b3\u30f3\u30d1\u30a4\u30eb\u3067\u304d\u306a\u3044\u306e\u306f\u3061\u3087\u3063\u3068\u3081\u3093\u3069\u304f\u3055\u3044\u3002
      \n\u4ed5\u7d44\u307f\u4e0a\u4ed5\u65b9\u306a\u3044\u3093\u3060\u3051\u3069\u3002<\/p>\n

      \u4e0d\u601d\u8b70\u306a\u306e\u306f\u3001\u304a\u308c\u304a\u308c\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u306e\u65b9\u306fO0\u301cO3\u3067\u5b9f\u884c\u901f\u5ea6\u306b\u307b\u3068\u3093\u3069\u5909\u5316\u304c\u306a\u3044\u306e\u306b\u5bfe\u3057\u3001thrust::sort\u306e\u65b9\u306fO0\u3068O3\u306730\u500d\u3050\u3089\u3044\u901f\u5ea6\u304c\u9055\u3046\uff08O0\u3060\u30681.8\u79d2\u3050\u3089\u3044\u304b\u304b\u308b\uff09\u3053\u3068\u3002
      \n\u6c17\u304c\u5411\u3044\u305f\u3089thrust\u306e\u30bd\u30fc\u30b9\u8aad\u3093\u3067\u307f\u3088\u3046\u304b\u306a\u3002<\/p>\n

      \u4ee5\u4e0b\u30bd\u30fc\u30b9\u3002
      \n\u304a\u308c\u304a\u308c\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u306e\u30ab\u30fc\u30cd\u30eb<\/p>\n

      __global__ void bitonic_sort(int *data, int data_size, int chunk_size, int sub_chunk_size)\r\n{\r\n    int thread_index = blockDim.x * blockIdx.x + threadIdx.x;\r\n\r\n    if (thread_index < data_size \/ 2) {\r\n        int half_chunk_size = chunk_size \/ 2;\r\n        int chunk_index = thread_index \/ half_chunk_size;\r\n        \r\n        int half_sub_chunk_size = sub_chunk_size \/ 2;\r\n        int sub_chunk_index = thread_index \/ half_sub_chunk_size;\r\n        \r\n        bool up = (chunk_index % 2 == 0);\r\n        int a = sub_chunk_size * sub_chunk_index + thread_index % half_sub_chunk_size;\r\n        int b = a + half_sub_chunk_size;\r\n        \r\n        int va = data[a];\r\n        int vb = data[b];\r\n        if (va > vb == up) {\r\n            data[a] = vb;\r\n            data[b] = va;\r\n        }\r\n    }\r\n}<\/pre>\n
      thrust::sort \u3092\u4f7f\u3063\u305f\u30b3\u30fc\u30c9\u306e\u4e3b\u8981\u90e8\u5206<\/p>\n
      #include <iostream>\r\n#include <cstdlib>\r\n#include "Timer.hpp"\r\n#include <thrust\/host_vector.h>\r\n#include <thrust\/device_vector.h>\r\n#include <thrust\/copy.h>\r\n#include <thrust\/sort.h>\r\n\r\nvoid thrust_sort(int n, int randomseed)\r\n{\r\n    std::cout << "Start thrust_sort\\n";\r\n    \r\n    \/\/ \u6642\u9593\u8a08\u6e2c\u30bf\u30a4\u30de\u30fc\r\n    Timer include_memcpy_timer, kernel_timer;\r\n    \r\n    \/\/ \u5143\u30c7\u30fc\u30bf\u306e\u4f5c\u6210\r\n    thrust::host_vector<int> hData;\r\n    hData.reserve(n);\r\n    srand(randomseed);\r\n    for (int i = 0; i < n; i++) {\r\n        double rnd = (double)rand() \/ RAND_MAX;\r\n        hData.push_back(rnd * n);\r\n    }\r\n    \r\n    \/\/ \u30db\u30b9\u30c8\u304b\u3089\u30c7\u30d0\u30a4\u30b9\u3078\u30b3\u30d4\u30fc\r\n    cudaDeviceSynchronize();\r\n    include_memcpy_timer.Start();\r\n    thrust::host_vector<int> dData(n);\r\n    thrust::copy(hData.begin(), hData.end(), dData.begin());\r\n    \r\n    \/\/ \u30bd\u30fc\u30c8\r\n    cudaDeviceSynchronize();\r\n    kernel_timer.Start();\r\n    thrust::sort(dData.begin(), dData.end());\r\n    cudaDeviceSynchronize();\r\n    kernel_timer.End();\r\n    \r\n    \/\/ \u30c7\u30d0\u30a4\u30b9\u304b\u3089\u30db\u30b9\u30c8\u3078\u30b3\u30d4\u30fc\r\n    thrust::copy(dData.begin(), dData.end(), hData.begin());\r\n    cudaDeviceSynchronize();\r\n    include_memcpy_timer.End();\r\n    \r\n    \/\/ \u7d50\u679c\u306e\u8868\u793a\r\n    std::cout << "Time: " << kernel_timer.GetSeconds() << " sec. (include memcpy: " << include_memcpy_timer.GetSeconds() << " sec.)\\n";\r\n    std::cout << "Result: " << n << " elements\\n";\r\n    for (int i = 0; i < ((n < 128)?n:128); i++) {\r\n        std::cout << hData[i] << ", ";\r\n    }\r\n    std::cout << "\\nDone.\\n";\r\n}<\/pre>\n","protected":false},"excerpt":{"rendered":"
      \u7d04100\u4e07\u8981\u7d20\uff081024 * 1024)\u306e32bit\u6574\u6570\u306e\u30bd\u30fc\u30c8\u901f\u5ea6\u3092\u4ee5\u4e0b\u306e\u30d7\u30ed\u30b0\u30e9\u30e0\u3067\u6e2c\u3063\u3066\u307f\u305f\u3002
      \n\u5168\u3066 C++\u3067\u66f8\u304d\u300164bit\u3001O3\u3067\u30b3\u30f3\u30d1\u30a4\u30eb\u3057\u305f\u3002
      \n\u5b9f\u884c\u3057\u305f\u30de\u30b7\u30f3\u306fCore 2 Duo 2GHz, Geforce 9400M\u3002<\/p>\n
        \n
      1. std::sort\u3067\u30bd\u30fc\u30c8\u3059\u308b\u30d7\u30ed\u30b0\u30e9\u30e0\u3002GPU\u975e\u4f7f\u7528\u3002\u3002\u30b3\u30f3\u30d1\u30a4\u30e9: clang++\u3002<\/li>\n
      2. \u524d\u56dePython\u3067\u66f8\u3044\u305f\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8<\/a>\u3092CUDA(Driver API)\u306b\u305d\u306e\u307e\u307e\u79fb\u690d\u3057\u305f\u3082\u306e\u3002\u30db\u30b9\u30c8\u30b3\u30fc\u30c9\u30b3\u30f3\u30d1\u30a4\u30e9: clang++\u3001\u30ab\u30fc\u30cd\u30eb\u30b3\u30f3\u30d1\u30a4\u30e9: nvcc\u3002<\/li>\n
      3. thrust::sort\u3067\u30bd\u30fc\u30c8\u3059\u308b\u30d7\u30ed\u30b0\u30e9\u30e0\u3002CUDA Runtime API\u4f7f\u7528\u3002\u30b3\u30f3\u30d1\u30a4\u30e9: nvcc\u3002<\/li>\n<\/ol>\n
        \uff08\u30bd\u30fc\u30b9\u306f\u672c\u6587\u306e\u6700\u5f8c\u306b\u8f09\u305b\u308b\u304c\u3001\u5168\u90e8\u306e\u305b\u308b\u3068\u9577\u304f\u306a\u308b\u306e\u3067\u3001\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u306e\u30ab\u30fc\u30cd\u30eb\u3068\u3001thrust::sort\u3092\u4f7f\u3063\u305f\u30d7\u30ed\u30b0\u30e9\u30e0\u306e\u4e3b\u8981\u306a\u90e8\u5206\u306e\u307f\u306e\u305b\u308b\u3002\uff09<\/p>\n
        \u7d50\u679c\u306f\u4ee5\u4e0b\u306e\u3068\u304a\u308a\u3002<\/p>\n
          \n
        1. std::sort \u7d040.37\u79d2<\/li>\n
        2. \u304a\u308c\u304a\u308c\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8 \u7d041.08\u79d2\uff08\u3046\u30610.01\u79d2\u306fCPU-GPU\u9593\u306e\u30e1\u30e2\u30ea\u30b3\u30d4\u30fc\uff09<\/li>\n
        3. thrust::sort \u7d040.05\u79d2\uff08\u3046\u30610.01\u79d2\u306fCPU-GPU\u9593\u306e\u30e1\u30e2\u30ea\u30b3\u30d4\u30fc\uff09<\/li>\n<\/ol>\n
          \u3055\u3059\u304c\u306bC++\u306estd::sort\u306f\u3001Python\u306esorted\u3088\u308a\u304b\u306a\u308a\u901f\u3044\u3002
          \n\u304a\u308c\u304a\u308c\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u306fGPU\u5411\u3051\u306e\u6700\u9069\u5316\u3092\u4e00\u5207\u3057\u3066\u306a\u3044\u306e\u3067\u9045\u3044\u306e\u306f\u4ed5\u65b9\u306a\u3044\u90e8\u5206\u3082\u3042\u308b\u3093\u3060\u3051\u3069\u3001CPU\u306e\u307f\u3067\u5b9f\u884c\u3059\u308bstd::sort\u3088\u308a\u9045\u3044\u306e\u306f\u3061\u3087\u3063\u3068\u3073\u3063\u304f\u308a(*_*;
          \nGPU\u304c\u53e4\u304f\u3066\u30e1\u30e2\u30ea\u30a2\u30af\u30bb\u30b9\u306b\u5bfe\u3059\u308b\u30b3\u30a2\u30ec\u30c3\u30b7\u30f3\u30b0\u304c\u52b9\u304d\u3065\u3089\u3044\u304b\u3089\u304b\u306a\u3041\u3002
          \n\u3067\u3001thrust::sort\u8d85\u901f\u3044\u3002
          \n\u3057\u304b\u3057thrust by CUDA\u306fnvcc\u3058\u3083\u306a\u3044\u3068\u307e\u3068\u3082\u306b\u30b3\u30f3\u30d1\u30a4\u30eb\u3067\u304d\u306a\u3044\u306e\u306f\u3061\u3087\u3063\u3068\u3081\u3093\u3069\u304f\u3055\u3044\u3002
          \n\u4ed5\u7d44\u307f\u4e0a\u4ed5\u65b9\u306a\u3044\u3093\u3060\u3051\u3069\u3002<\/p>\n
          \u4e0d\u601d\u8b70\u306a\u306e\u306f\u3001\u304a\u308c\u304a\u308c\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u306e\u65b9\u306fO0\u301cO3\u3067\u5b9f\u884c\u901f\u5ea6\u306b\u307b\u3068\u3093\u3069\u5909\u5316\u304c\u306a\u3044\u306e\u306b\u5bfe\u3057\u3001thrust::sort\u306e\u65b9\u306fO0\u3068O3\u306730\u500d\u3050\u3089\u3044\u901f\u5ea6\u304c\u9055\u3046\uff08O0\u3060\u30681.8\u79d2\u3050\u3089\u3044\u304b\u304b\u308b\uff09\u3053\u3068\u3002
          \n\u6c17\u304c\u5411\u3044\u305f\u3089thrust\u306e\u30bd\u30fc\u30b9\u8aad\u3093\u3067\u307f\u3088\u3046\u304b\u306a\u3002<\/p>\n
          \u4ee5\u4e0b\u30bd\u30fc\u30b9\u3002
          \n\u304a\u308c\u304a\u308c\u30d0\u30a4\u30c8\u30cb\u30c3\u30af\u30bd\u30fc\u30c8\u306e\u30ab\u30fc\u30cd\u30eb<\/p>\n
          __global__ void bitonic_sort(int *data, int data_size, int chunk_size, int sub_chunk_size)\r\n{\r\n    int thread_index = blockDim.x * blockIdx.x + threadIdx.x;\r\n\r\n    if (thread_index < data_size \/ 2) {\r\n        int half_chunk_size = chunk_size \/ 2;\r\n        int chunk_index = thread_index \/ half_chunk_size;\r\n        \r\n        int half_sub_chunk_size = sub_chunk_size \/ 2;\r\n        int sub_chunk_index = thread_index \/ half_sub_chunk_size;\r\n        \r\n        bool up = (chunk_index % 2 == 0);\r\n        int a = sub_chunk_size * sub_chunk_index + thread_index % half_sub_chunk_size;\r\n        int b = a + half_sub_chunk_size;\r\n        \r\n        int va = data[a];\r\n        int vb = data[b];\r\n        if (va > vb == up) {\r\n            data[a] = vb;\r\n            data[b] = va;\r\n        }\r\n    }\r\n}<\/pre>\n
          thrust::sort \u3092\u4f7f\u3063\u305f\u30b3\u30fc\u30c9\u306e\u4e3b\u8981\u90e8\u5206<\/p>\n
          #include <iostream>\r\n#include <cstdlib>\r\n#include "Timer.hpp"\r\n#include <thrust\/host_vector.h>\r\n#include <thrust\/device_vector.h>\r\n#include <thrust\/copy.h>\r\n#include <thrust\/sort.h>\r\n\r\nvoid thrust_sort(int n, int randomseed)\r\n{\r\n    std::cout << "Start thrust_sort\\n";\r\n    \r\n    \/\/ \u6642\u9593\u8a08\u6e2c\u30bf\u30a4\u30de\u30fc\r\n    Timer include_memcpy_timer, kernel_timer;\r\n    \r\n    \/\/ \u5143\u30c7\u30fc\u30bf\u306e\u4f5c\u6210\r\n    thrust::host_vector<int> hData;\r\n    hData.reserve(n);\r\n    srand(randomseed);\r\n    for (int i = 0; i < n; i++) {\r\n        double rnd = (double)rand() \/ RAND_MAX;\r\n        hData.push_back(rnd * n);\r\n    }\r\n    \r\n    \/\/ \u30db\u30b9\u30c8\u304b\u3089\u30c7\u30d0\u30a4\u30b9\u3078\u30b3\u30d4\u30fc\r\n    cudaDeviceSynchronize();\r\n    include_memcpy_timer.Start();\r\n    thrust::host_vector<int> dData(n);\r\n    thrust::copy(hData.begin(), hData.end(), dData.begin());\r\n    \r\n    \/\/ \u30bd\u30fc\u30c8\r\n    cudaDeviceSynchronize();\r\n    kernel_timer.Start();\r\n    thrust::sort(dData.begin(), dData.end());\r\n    cudaDeviceSynchronize();\r\n    kernel_timer.End();\r\n    \r\n    \/\/ \u30c7\u30d0\u30a4\u30b9\u304b\u3089\u30db\u30b9\u30c8\u3078\u30b3\u30d4\u30fc\r\n    thrust::copy(dData.begin(), dData.end(), hData.begin());\r\n    cudaDeviceSynchronize();\r\n    include_memcpy_timer.End();\r\n    \r\n    \/\/ \u7d50\u679c\u306e\u8868\u793a\r\n    std::cout << "Time: " << kernel_timer.GetSeconds() << " sec. (include memcpy: " << include_memcpy_timer.GetSeconds() << " sec.)\\n";\r\n    std::cout << "Result: " << n << " elements\\n";\r\n    for (int i = 0; i < ((n < 128)?n:128); i++) {\r\n        std::cout << hData[i] << ", ";\r\n    }\r\n    std::cout << "\\nDone.\\n";\r\n}<\/pre>\n","protected":false},"author":1,"featured_media":0,"comment_status":"open","ping_status":"open","sticky":false,"template":"","format":"standard","meta":{"footnotes":""},"categories":[32],"tags":[458,437,289,459],"class_list":["post-2678","post","type-post","status-publish","format-standard","hentry","category-tech","tag-458","tag-c","tag-cuda","tag-thrust"],"_links":{"self":[{"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/posts\/2678","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/comments?post=2678"}],"version-history":[{"count":0,"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/posts\/2678\/revisions"}],"wp:attachment":[{"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/media?parent=2678"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/categories?post=2678"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/peta.okechan.net\/blog\/wp-json\/wp\/v2\/tags?post=2678"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}