CUDAによる再帰関数の高速化は可能なのか

タイトルのとおりなのですが、再帰関数をCUDAを使って高速化するというのは可能なのでしょうか？
例えば、配列の各要素に対してある操作をするというのはCUDAを使うことで並列に処理できて高速化が実現されるというのはよくみます。再帰関数の場合には、処理をどう各スレッドに分ければ良いのかわかりませんでした。

試しにフィボナッチ数列のn番目の数字を求めるコードを書いてみました。

cuda
1#include <stdio.h>
2
3__device__ void fib_cuda(int n, int *out) {
4  if (n < 3) {
5    *out = 1;
6  } else {
7    int *p, *pp;
8    p = (int*)malloc(sizeof(int));
9    pp = (int*)malloc(sizeof(int));
10
11    fib_cuda(n - 1, p);
12    fib_cuda(n - 2, pp);
13
14    *out = *p + *pp;
15    free(p);
16    free(pp);
17  }
18}
19
20__global__ void fib(int n, int *out) {
21  fib_cuda(n, out);
22}
23
24int main() {
25  int *out;
26  int *d_out;
27  out = (int*)malloc(sizeof(int));
28
29  cudaMalloc((void**)&d_out, sizeof(int));
30  cudaMemcpy(d_out, out, sizeof(int), cudaMemcpyHostToDevice);
31
32  fib<<<1, 1>>>(20, d_out);
33  cudaMemcpy(out, d_out, sizeof(int), cudaMemcpyDeviceToHost);
34  printf("%d\n", *out);
35
36  cudaFree(d_out);
37  free(out);
38}
39

このコードは動きはするのですが、単にGPU上でCPUと同じ処理をしただけでGPUの恩恵は皆無です。もし再帰関数の高速化も可能ならば上のコードはどう書かれるべきなのでしょうか？

抽象的な質問で申し訳ないのですがよろしくおねがいします。

fana

2021/12/15 02:28

関係あるかどうかわかりませんが「フィボナッチ数列 gpu」でググると，検索結果の先頭らへんに，文献： ”CUDA を用いたメッセージ送信型並行計算 Actor モデルの実装”, 高柳, 脇田, IPSJ SIG Technial Report Vol.2012-HPC-135 No.37, 2012/8/3 なるPDFがヒットしました．（私は内容を読んではいませんが）実験の例題としてフィボナッチ数列を扱っているようです．何かしら並列性のある演算を実装している話かもしれません．（図9を見るに「CPUの方が早いけどね」という話になっているようにも見えますが）

行動規範の内容に同意します

回答2件

ベストアンサー

cudaは、nVIDIAのGPUを使ってSIMD命令(Single Instruction Multiple Data命令)による機械命令レベルのもっとも細かい並列化を行うためのライブラリです。
つまり、16個なり64個なりのデータを対象にして、ひとつの加算なり乗算なりの命令を出すことで同時に6個なり64個なりの計算を実行することによって高速化を行うためのものです。

一般的には再帰関数はSIMD命令に変換することはできませんが、tail recursionと呼ばれる再帰関数については自動最適化コンパイラがループに変換してSIMD命令を出力できる場合があります。
インテルチップやARMチップに内蔵されているSIMD命令についてはこれを行っているコンパイラは存在するはずです。
しかし、cudaのように決め打ちでライブラリを呼び出すような方式ですと自動最適化コンパイラが介在する余地がありませんので、現在はできませんし、将来的にも難しいと思います。

また、例に挙げられているようなフィボナッチ数列タイプの再帰関数では手動でもループに変換することが困難ですので、将来的にも無理でしょう。

投稿2021/12/04 13:18