numba for CUDA の使い方が分からない

Question

質問の背景

nunba for cuda を使用するとpythonでの計算が数百倍になるという記事を読んで、実際にやってみようと思いましたが、エラーが出てしまいます。

1万まで1を足していくプログラムを試しに作ってみましたが、

from numba import cuda

@cuda.jit
def hello():
    count = 0
    for i in range(10000):
        count+=1
    return count
    
print(hello())

ValueError: 
Kernel launch configuration was not specified. Use the syntax:

kernel_function[blockspergrid, threadsperblock](arg0, arg1, ..., argn)

See https://numba.pydata.org/numba-doc/latest/cuda/kernels.html#kernel-invocation for help.

このようなエラーが出てきてしまいました。

おそらく使い方が間違っていると思いますが、ネットに情報が少なくどうすればいいかよく分かりません。

###質問

例えば質問の背景に示したように０から1を１万回足すプログラムをnumba for cudaを用いて高速化したい場合、どのように記述すればいいですか？

Comments

[8524ba23]

参考にした記事、調査、試したことの詳細を記載ください。
なお、質問への追記・修正の依頼や回答がついていいるにもかかわらず、それらに反応がない過去の質問が多いようですが、受け答えができない（あなたが放置する）ユーザーとみなされると（第三者にも）放置される可能性が高いことを認識ください。

Answer 1

1を１万回足すプログラム

を、「1万個の1を足すプログラム」と置き換えて、以下書きます

ただし、一つの変数に結果を合算するのは、
CUDAで配列の総和を求めてみた
等に書かれてるように、CUDAで速度を出すために並列化するにはいろいろと工夫が要るケースで、この回答のCUDAを用いたコードはそこを頑張ってないので、

import numpy as np
import time

number_of_maximum_loop = 10000

start = time.time()
count = np.sum(np.ones(number_of_maximum_loop, dtype="int"))
elapsed_time = time.time() - start
print(count)
print("elapsed_time:{0}".format(elapsed_time) + "[sec]")

に速度で圧倒的に負けますので、

numba for cudaを用いて高速化

は満たしてません
とりあえず、こう書けばエラーが出なくなります、というだけです
(注：【追記】に書いたように、データのサイズを大きくしたらCUDAとnumpyの速度が逆転しました)

　
numba cuda does not produce correct result with += (gpu reduction needed?)
を参考にした、質問に掲載のエラーを回避したコード

from numba import cuda
import time

@cuda.jit
def cuda_kernel(number_of_maximum_loop, result):
    i = cuda.grid(1)
    if i < number_of_maximum_loop:
        cuda.atomic.add(result, 0, 1)

number_of_maximum_loop = 10000
threadsperblock = 1024
blockspergrid = (number_of_maximum_loop + (threadsperblock - 1)) // threadsperblock

result = cuda.to_device([0])
cuda_kernel[blockspergrid, threadsperblock](number_of_maximum_loop, result)
count = result.copy_to_host()[0]
# 1回目はコンパイルするため2回目を測定
start = time.time()
result = cuda.to_device([0])
cuda_kernel[blockspergrid, threadsperblock](number_of_maximum_loop, result)
count = result.copy_to_host()[0]
elapsed_time = time.time() - start
print(count)
print("elapsed_time:{0}".format(elapsed_time) + "[sec]")

　
.
GPU Reduction
を参考にしたコード

import numpy as np
from numba import cuda
import time

@cuda.reduce
def sum_reduce(a, b):
    return a + b

number_of_maximum_loop = 10000

count = sum_reduce(np.ones(number_of_maximum_loop, dtype="int"))
# 1回目はコンパイルするため2回目を測定
start = time.time()
count = sum_reduce(np.ones(number_of_maximum_loop, dtype="int"))
elapsed_time = time.time() - start
print(count)
print("elapsed_time:{0}".format(elapsed_time) + "[sec]")

　
.
【追記】
Google ColabでGPUを有効にして確認したところ、

number_of_maximum_loop = 10000

を

number_of_maximum_loop = 1000000000

に変えたら、上記の一つ目のCUDAを用いたコードの速度はnumpyだけのコードの速度をやや上回りましたので、扱うデータのサイズが大きくなるとCUDAが有利になるようです
CUDAのコードを、回答の最初に提示したWebページに書かれてるような合算の高速化に対応させたものにすれば、さらにCUDAの方が有利になると思います (未確認)
ただし、「number_of_maximum_loop = 1000000000」の場合でも上記二つ目のCUDAを用いたコードはnumpyだけのコードよりもやや遅いです
もしかしたら、下記で指摘されてるバグが、まだ直ってないのかもしれません
Numba CUDA vectorize and reduce decorators slower than expected #2266

Answer 2

エラーメッセージの最後に表示されている文言にある URL にアクセスしてみると、

See https://numba.pydata.org/numba-doc/latest/cuda/kernels.html#kernel-invocation for help.

以下の様に書かれています。

Kernel invocation

A kernel is typically launched in the following way:

threadsperblock = 32
blockspergrid = (an_array.size + (threadsperblock - 1)) // threadsperblock
increment_by_one[blockspergrid, threadsperblock](an_array)

なので、hello() 関数の呼び出しは以下の様に記述するのではないでしょうか([10, 20] は適当な値です)。

print(hello[10, 20]())