CUDA Thrustライブラリを用いて、CUDAで配列の総和を求める方法

本記事では、CUDAのThrustのリダクション（reduction）計算の機能を用いて、GPU側で配列の総和を求める方法を紹介します。

リダクションとは、配列の要素（多数）を一つの結果にまとめることで、例えば配列の全ての要素の総和や最大値、最小値などを求める処理が該当します。

CUDA Thrustについて

ThrustはCUDA C++のテンプレートライブラリで、CUDAをインストールすると自動的に付いてくるライブラリです。

これはC++のSTL（Standard Template Library）と呼ばれるC++の開発でよく使われる標準ライブラリと似た文法で書けることが特徴で、通常のCUDAのプログラミングと比較すると、CPUからGPU側へのデータの転送などを簡易な文法で実現することができます。

よってThrustを使うことで、C++に慣れている人にとっては比較的少ない労力でCUDAを用いたGPUプログラミングをすることができます。

一方、実装されている処理には限りがあり、何でもできるわけではありません。Thrustに実装されている機能の例としては、配列の総和を求めたり、配列をソートしたりする機能が挙げられます。

CUDA Thrustを用いて配列の要素の総和を求めるプログラム

以下に、GPU側で配列の総和と求めるプログラムを紹介します。

#include <iostream>
#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/reduce.h>

int main()
{
    // Creating array in host
    const int N = 100;
    int* data_host;
    data_host = (int*)malloc(sizeof(int) * N);
    for (int i = 0; i < N; i++) {
        data_host[i] = i + 1;
    }

    // Data transfer (CPU -> GPU) 
    int* data_dev;
    cudaMalloc(&data_dev, sizeof(int) * N);
    cudaMemcpy(data_dev, data_host, sizeof(int) * N, cudaMemcpyHostToDevice);

    // Calculate the sum in GPU by thrust
    thrust::device_vector<int> d_vec(data_dev, data_dev + N);
    int sum1 = thrust::reduce(d_vec.begin(), d_vec.end());
    int sum2 = thrust::reduce(d_vec.begin(), d_vec.end(), (int)100);   // ininial value = 100;

    // Confirmation
    std::cout << "sum1 = " << sum1 << std::endl;
    std::cout << "sum2 = " << sum2 << std::endl;
   
    return 0;
}

このreduction機能を使うためには、「thrust/reduce.h」をインクルードする必要があります。

8～14行目でint型の100個の要素からなる配列を用意し、1～100の値を入力します。その後、16～19行目でデータをCPU側からGPU側に転送します。

今回のポイントとなる総和を求める部分は、23行目と24行目です。

23行目「int sum1 = thrust::reduce(d_vec.begin(), d_vec.end());」は最も簡単な処理で、配列の総和を算出します。

24行目「int sum2 = thrust::reduce(d_vec.begin(), d_vec.end(), (int)100);」は、3つ目の引数として「100」を与えていますが、これは総和を求める際の初期値として100を与えることができるそうです。

ちなみに、今回計算する1～100の総和は5050となります。

実行結果は以下の通りです。正常に総和を計算できていることがわかります。