Remarks
이 글은 numba.cuda를 기반으로 작성되었습니다.
Parallel computing에서 kernel 함수 내부의 불필요한 if문은 속도 저하의 원인 중 하나입니다.
만약 프로그램이 실행되는 동안 바뀌지 않는 조건을 포함하는 if문이 존재한다면, closure를 사용하여 동일한 조건을 확인하는 연산들을 한 번으로 줄일 수 있습니다.
그리고 특정한 data type만을 인식할 수 있는 kernel(device) function의 제약을 깔끔하게 처리해줄 수 있습니다. 가령, str type을 가진 변수에 대한 조건을 만들어 줄 수 있죠. 저는 보통 전역 class로 parameter를 관리하고 있는데 class 혹은 dictionary 등의 data의 값을 읽어올 수 있는 것이 굉장히 편리합니다.
다음은 data의 각 행을 하나의 thread에 할당하여 각 행의 값들을 집계(aggregation)하는 예제입니다.
집계하는 방식을 class G의 METHOD attribute로 정하고, 만약 METHOD가 select라면 G.IDX_SELECT의 값을 반환하는 kernel function을 반환합니다.
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
from numba import cuda
def get_COMPUTE_AGGREGATION(G):
METHOD:str = G.METHOD
@cuda.jit
def COMPUTE_AGGREGATION(rst_vec, data_mat):
tid = cuda.grid(1)
rst_vec[tid] = DEV_GET_AGGREGATION(data_mat[tid])
if METHOD is 'sum':
@cuda.jit(device=True)
def DEV_GET_AGGREGATION(data_vec):
rst = 0
for data in range(data_vec):
rst += data
return rst
elif METHOD is 'select':
IDX_SELECT = G.IDX_SELECT
@cuda.jit(device=True)
def DEV_GET_AGGREGATION(data_vec):
rst = data_vec[IDX_SELECT]
return rst
else:
raise ValueError
return COMPUTE_AGGREGATION
1
2
3
4
5
6
7
8
9
10
11
12
13
class G:
METHOD = 'select'
IDX_SELECT = 2
COMPUTE_AGGREGATION = get_COMPUTE_AGGREGATION(G)
rst_vec = np.zeros(5)
data_mat = np.stack([range(10) for _ in range(5)])
COMPUTE_AGGREGATION[1, len(rst_vec)](rst_vec, data_mat)
print(data_mat, '\n')
print(rst_vec)
[[0 1 2 3 4 5 6 7 8 9]
[0 1 2 3 4 5 6 7 8 9]
[0 1 2 3 4 5 6 7 8 9]
[0 1 2 3 4 5 6 7 8 9]
[0 1 2 3 4 5 6 7 8 9]]
[2. 2. 2. 2. 2.]
PREVIOUSEtc