소개

Jetson Nano의 사용법으로 시뮬레이션을 잊을 수는 없습니다. JetPack에 사전 설치된 CUDA 툴킷에는 시뮬레이션 관련 샘플 프로그램이 포함되어 있으며 예를 들어 N-Body 시뮬레이션이 유명합니다. 이 샘플 프로그램은 물론 Jetson Nano에서도 작동합니다. 저도 처음에는 이 N-Body 시뮬레이션 프로그램을 개조해 여러가지 시도하려고 했습니다만, 이 프로그램은 꽤 복잡해서, 조금 시험하기 위한 편지지로서 사용하기는 어렵다고 판단했습니다.
다음에 생각한 것이, CUDA 로 실장한 라이프 게임이 전세계에서 여러가지 공개되고 있을 것이라고 생각했으므로, 그것을 히나형으로 하는 것입니다. 굉장히 완성도가 높은 실장예 "Jetson Nano로 GPGPU (CUDA) 프로그래밍 ~라이프 게임 개발~" 를 발견했습니다만, 약간의 실험의 편지지에 사용하기에는 너무 정교한 실장이라고 하는 느낌입니다. 한층 더 조사하면 "Python에 Curses 라이브러리가 있었기 때문에 간단한 라이프 게임을 만들어 보는 테스트" 라고 하는 기사가 있었으므로 이것을 참고로 하겠고, 표시 부분은 Python 의 Curses 라이브러리를 사용하는 것으로 하고, CUDA 부분은 PyCUDA 를 사용해 스스로 CUDA 프로그래밍 했습니다.
그렇다고 해서, 시뮬레이션을 시험하는 히나형으로서 자작한 간단한 코드를 여기서 소개하겠습니다.

구현

작성한 코드는 GitHub tsutof/game_of_life_pycuda에서 게시 중입니다.

사용 목적이 CUDA에 의한 시뮬레이션의 편지지이므로, 간결함을 유의했습니다.

Python 프로그램이지만 인접 셀을 확인하고 상태 전이하는 계산 부분은 일반 CUDA C 언어 코드입니다. 이 코드는 런타임에 컴파일되므로 프로그램 실행 시작 후 처음에 이 코드가 CUDA 코어에서 실행될 때는 호흡을 기다리는 느낌입니다. 물론, 그 후에는 빠르게 실행됩니다. 보시다시피 알 수 있도록이 코드는 하나의 셀만 대상으로 하고 있습니다. NVIDIA GPU가 내장한 많은 CUDA 코어에서 이 코드는 병렬로 실행됩니다.

lifegame.py

def get_next_state_gpu(state, next_state):
    height, width = state.shape

    mod = SourceModule("""
        __global__ void get_next_state(int *state, int *nextState, int height, int width)
        {
            unsigned int ix = threadIdx.x + blockIdx.x * blockDim.x;
            unsigned int iy = threadIdx.y + blockIdx.y * blockDim.y;
            unsigned int idx = iy * width + ix;
            int sum = 0;
            int val, nextVal;
            if (ix >= width || iy >= height) {
                return;
            }
            val = state[idx];
            sum += state[((iy - 1) % height) * width + ((ix - 1) % width)];
            sum += state[((iy - 1) % height) * width + (ix       % width)];
            sum += state[((iy - 1) % height) * width + ((ix + 1) % width)];
            sum += state[(iy       % height) * width + ((ix - 1) % width)];
            sum += state[(iy       % height) * width + ((ix + 1) % width)];
            sum += state[((iy + 1) % height) * width + ((ix - 1) % width)];
            sum += state[((iy + 1) % height) * width + (ix       % width)];
            sum += state[((iy + 1) % height) * width + ((ix + 1) % width)];
            if (val == 0 && sum == 3) {
                nextVal = 1;
            }
            else if (val != 0 && (sum >= 2 && sum <= 3)) {
                nextVal = 1;
            }
            else {
                nextVal = 0;
            }
            nextState[idx] = nextVal;
        }
        """)
    kernel_func = mod.get_function("get_next_state")

    blk_dim = (32, 32, 1)
    grid_dim = ((width + blk_dim[0] - 1) // blk_dim[0], \
        (height + blk_dim[1] - 1) // blk_dim[1], 1)
    kernel_func(
        drv.In(state), drv.Out(next_state), \
        np.int32(height), np.int32(width), \
        block=blk_dim, grid=grid_dim)

요약

이 프로그램에는 CPU에서 모두 처리하는 모드도 추가했습니다. 명령행 옵션으로 전환할 수 있습니다. 48x157 = 7536 셀의 경우 CUDA 코어 처리에서 약 100배 가속화되었습니다. (인접 셀을 체크하여 상태 천이하는 계산 부분만으로, 표시 부분은 제외한 배율입니다. Jeton Nano MAXN 모드) CPU만으로 계산하는 경우는, 싱글 thread 처리이므로, 이것을 멀티 thread화하면, CPU 처리의 경우도 몇 배나 고속화된다고 생각합니다만, 그래도 CUDA 코어에 의한 처리는 CPU 처리에 비해 압도적으로 고속이라고 하는 결과가 되었습니다.

이상입니다.