teratail
回答率
85.35%
teratail header banner
teratail header banner
質問するログイン新規登録
トップCUDAに関する質問numba for CUDA の使い方が分からない

編集履歴

回答編集履歴

2

説明の追記

2021/11/01 05:19

投稿

jbpb0
jbpb0

スコア7653

test CHANGED
@@ -20,9 +20,13 @@
20
20
 
21
21
 
22
22
 
23
+ number_of_maximum_loop = 10000
24
+
25
+
26
+
23
27
  start = time.time()
24
28
 
25
- count = np.sum(np.ones(10000, dtype="int"))
29
+ count = np.sum(np.ones(number_of_maximum_loop, dtype="int"))
26
30
 
27
31
  elapsed_time = time.time() - start
28
32
 
@@ -43,6 +47,8 @@
43
47
  は満たしてません
44
48
 
45
49
  とりあえず、こう書けばエラーが出なくなります、というだけです
50
+
51
+ (注：【追記】に書いたように、データのサイズを大きくしたらCUDAとnumpyの速度が逆転しました)
46
52
 
47
53
 
48
54
 
@@ -80,6 +86,14 @@
80
86
 
81
87
 
82
88
 
89
+ result = cuda.to_device([0])
90
+
91
+ cuda_kernel[blockspergrid, threadsperblock](number_of_maximum_loop, result)
92
+
93
+ count = result.copy_to_host()[0]
94
+
95
+ # 1回目はコンパイルするため2回目を測定
96
+
83
97
  start = time.time()
84
98
 
85
99
  result = cuda.to_device([0])
@@ -108,6 +122,8 @@
108
122
 
109
123
  ```python
110
124
 
125
+
126
+
111
127
  import numpy as np
112
128
 
113
129
  from numba import cuda
@@ -124,9 +140,17 @@
124
140
 
125
141
 
126
142
 
143
+ number_of_maximum_loop = 10000
144
+
145
+
146
+
147
+ count = sum_reduce(np.ones(number_of_maximum_loop, dtype="int"))
148
+
149
+ # 1回目はコンパイルするため2回目を測定
150
+
127
151
  start = time.time()
128
152
 
129
- count = sum_reduce(np.ones(10000, dtype="int"))
153
+ count = sum_reduce(np.ones(number_of_maximum_loop, dtype="int"))
130
154
 
131
155
  elapsed_time = time.time() - start
132
156
 
@@ -135,3 +159,37 @@
135
159
  print("elapsed_time:{0}".format(elapsed_time) + "[sec]")
136
160
 
137
161
  ```
162
+
163
+
164
+
165
+ 　
166
+
167
+ 　.
168
+
169
+ 【追記】
170
+
171
+ Google ColabでGPUを有効にして確認したところ、
172
+
173
+ ```python
174
+
175
+ number_of_maximum_loop = 10000
176
+
177
+ ```
178
+
179
+ を
180
+
181
+ ```python
182
+
183
+ number_of_maximum_loop = 1000000000
184
+
185
+ ```
186
+
187
+ に変えたら、上記の一つ目のCUDAを用いたコードの速度はnumpyだけのコードの速度をやや上回りましたので、扱うデータのサイズが大きくなるとCUDAが有利になるようです
188
+
189
+ CUDAのコードを、回答の最初に提示したWebページに書かれてるような合算の高速化に対応させたものにすれば、さらにCUDAの方が有利になると思います (未確認)
190
+
191
+ ただし、「number_of_maximum_loop = 1000000000」の場合でも上記二つ目のCUDAを用いたコードはnumpyだけのコードよりもやや遅いです
192
+
193
+ もしかしたら、下記で指摘されてるバグが、まだ直ってないのかもしれません
194
+
195
+ [Numba CUDA vectorize and reduce decorators slower than expected #2266](https://github.com/numba/numba/issues/2266)

1

コード追加

2021/11/01 05:18

投稿

jbpb0
jbpb0

スコア7653

test CHANGED
@@ -2,13 +2,117 @@
2
2
 
3
3
 
4
4
 
5
- 1万個の1を足すプログラム
5
+ を、「1万個の1を足すプログラム」と置き換えて、以下書きます
6
+
7
+
8
+
9
+ ただし、一つの変数に結果を合算するのは、
10
+
11
+ [CUDAで配列の総和を求めてみた](https://qiita.com/gyu-don/items/ef8a128fa24f6bddd342)
12
+
13
+ 等に書かれてるように、CUDAで速度を出すために並列化するにはいろいろと工夫が要るケースで、この回答のCUDAを用いたコードはそこを頑張ってないので、
14
+
15
+ ```python
16
+
17
+ import numpy as np
18
+
19
+ import time
20
+
21
+
22
+
23
+ start = time.time()
24
+
25
+ count = np.sum(np.ones(10000, dtype="int"))
26
+
27
+ elapsed_time = time.time() - start
28
+
29
+ print(count)
30
+
31
+ print("elapsed_time:{0}".format(elapsed_time) + "[sec]")
32
+
33
+ ```
34
+
35
+ に速度で圧倒的に負けますので、
36
+
37
+
38
+
39
+ > numba for cudaを用いて高速化
40
+
41
+
42
+
43
+ は満たしてません
44
+
45
+ とりあえず、こう書けばエラーが出なくなります、というだけです
46
+
47
+
48
+
49
+ 　
50
+
51
+ [numba cuda does not produce correct result with += (gpu reduction needed?)](https://stackoverflow.com/questions/53745701/numba-cuda-does-not-produce-correct-result-with-gpu-reduction-needed)
52
+
53
+ を参考にした、質問に掲載のエラーを回避したコード
54
+
55
+ ```python
56
+
57
+ from numba import cuda
58
+
59
+ import time
60
+
61
+
62
+
63
+ @cuda.jit
64
+
65
+ def cuda_kernel(number_of_maximum_loop, result):
66
+
67
+ i = cuda.grid(1)
68
+
69
+ if i < number_of_maximum_loop:
70
+
71
+ cuda.atomic.add(result, 0, 1)
72
+
73
+
74
+
75
+ number_of_maximum_loop = 10000
76
+
77
+ threadsperblock = 1024
78
+
79
+ blockspergrid = (number_of_maximum_loop + (threadsperblock - 1)) // threadsperblock
80
+
81
+
82
+
83
+ start = time.time()
84
+
85
+ result = cuda.to_device([0])
86
+
87
+ cuda_kernel[blockspergrid, threadsperblock](number_of_maximum_loop, result)
88
+
89
+ count = result.copy_to_host()[0]
90
+
91
+ elapsed_time = time.time() - start
92
+
93
+ print(count)
94
+
95
+ print("elapsed_time:{0}".format(elapsed_time) + "[sec]")
96
+
97
+ ```
98
+
99
+
100
+
101
+ 　
102
+
103
+ 　.
104
+
105
+ [GPU Reduction](https://numba.pydata.org/numba-doc/dev/cuda/reduction.html)
106
+
107
+ を参考にしたコード
6
108
 
7
109
  ```python
8
110
 
9
111
  import numpy as np
10
112
 
11
113
  from numba import cuda
114
+
115
+ import time
12
116
 
13
117
 
14
118
 
@@ -20,22 +124,14 @@
20
124
 
21
125
 
22
126
 
127
+ start = time.time()
128
+
23
- print(sum_reduce(np.ones((10000))))
129
+ count = sum_reduce(np.ones(10000, dtype="int"))
130
+
131
+ elapsed_time = time.time() - start
132
+
133
+ print(count)
134
+
135
+ print("elapsed_time:{0}".format(elapsed_time) + "[sec]")
24
136
 
25
137
  ```
26
-
27
- 参考
28
-
29
- [GPU Reduction](https://numba.pydata.org/numba-doc/dev/cuda/reduction.html)
30
-
31
-
32
-
33
- ただし、時間測ってないので、
34
-
35
-
36
-
37
- > numba for cudaを用いて高速化
38
-
39
-
40
-
41
- になってるかは不明