CUDAで配列の最大・最小値を算出するプログラム

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

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

CUDA Thrustについて

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

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

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

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

CUDAで配列の最大・最小値を算出するプログラム

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

#include <iostream>
#include <cuda_runtime.h>
#include <thrust/device_vector.h>
#include <thrust/functional.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 min_val = thrust::reduce(d_vec.begin(), d_vec.end(), INT_MAX, thrust::minimum<int>());
    int max_val = thrust::reduce(d_vec.begin(), d_vec.end(), INT_MIN, thrust::maximum<int>());

    // Confirmation
    std::cout << "min_val = " << min_val << std::endl;
    std::cout << "max_val = " << max_val << std::endl;

    return 0;
}

今回はCUDAのthrustライブラリのreduction機能を利用しており、これを利用して配列の最大値や最小値を求めるためには、「thrust/reduce.h」「thrust/functional.h」をインクルードする必要があります。

9～15行目でint型の100個の要素からなる配列を用意し、1～100の値を入力します。そのため、この配列の最大値は100であり、最小値は1です。

その後、17～20行目でデータをCPU側からGPU側に転送します。

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

24行目「int min_val = thrust::reduce(d_vec.begin(), d_vec.end(), INT_MAX, thrust::minimum());」は、INT型の最小値を算出します。第三引数に「INT_MAX」を与えていますが、これは最小値を求める過程での初期値であり、初期値に巨大な値を設定しています。

25行目「int max_val = thrust::reduce(d_vec.begin(), d_vec.end(), INT_MIN, thrust::maximum());」は、配列の最大値を求める処理です。

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