CUDA - 병렬 프로그래밍(스레드를 이용한 벡터의 합)

이번에는 앞전에 했던 블록을 이용한 벡터의 합을 스레드를 이용한 벡터의 합으로 구현해보려고 한다. 

사실 코드는 몇줄 바뀌지 않는다. 단순하게 블록에서 하던 일을 스레드에서 하게 바꾸는 작업만 하면 되는 일이니까...

앞에서 작성했던 코드와 비교해서 바뀐 부분은 빨간색 네모를 확인하면 된다.

그림 - add()에서 바뀐 부분

그림 - main()에서 바뀐 부분

앞에서 하드웨어의 제약으로 인해 한 번에 실행 가능한 블록의 수는 65,535개를 초과할 수 없다고 했었다. 동일하게, 하나의 커널을 실행하는 데 있어서 블록당 스레드 수 역시 하드웨어의 제약을 받는다. 구체적으로, 이 수는 우리가 3장에서 보았던 디바이스 속성 구조체의 'maxThreadsPerBlock' 멤버의 값을 초과해서는 안된다. 현재 이용가능한 많은 그래픽스 프로세서들에서 블록당 스레드의 최대 개수는 512개이다.

그렇다면 배열의 개수가 512개보다 더 클 때 어떻게 스레드 기반으로 두 벡터의 합을 수행할 수 있는지 생각해보자. 뭐 방법은 간단하다. 스레드와 블록을 같이 사용하는 방법이다.

그림 - 스레드와 블록을 같이 사용

여기서 'blockDim'이라는 내장 변수를 사용하는데 앞에서 설명했다시피, 각 블록당 스레드의 개수 이다. 여기서는 1차원의 블록을 사용하고 있기 때문에 그냥 'blockDim.x'만 사용했고, 2차원이라면 'y'도 사용 해야 할 것이다.

또 다른 변경은 커널 실행 자체에 있다. 실행할 N개의병렬 스레드들이 여전히 필요하지만, 우리에게 부과된 512개의 스레드 제한 수를 초과하지 않도록 다수의 블록들을 통해 스레드들을 발동시키길 원한다. 하나의 해결책은 임의로 블록 크기를 고정된 스레드의 개수로 지정하는 것이다. 예를 들어, 블록당 128개의 스레드들을 사용한다고 고정시켜보자.

그리고 블록을 생성하는 개수를 정할 차례인데, 여기서 묘안은 N/128이 정수의 나눗셈이라는 것이다. 만약 N이 127이라면 N/128은 정수로 0이 된다. 즉 0개의 블록을 생성하게 되므로 스레드도 실행되지 않게 된다. 하지만, 이렇게 구현해버리면, N이 정확히 128의 배수여야만 올바른 블록의 개수가 실행될 것이다. 이것은 좋지 못하다. 이를 해결하기 위한 일반적인 요령은 다음과 같다.

그림 - 커널의 다른 실행 방법

여기서 N이 10개 이므로, 블록은 10개 스레드는 128개가 된다. 즉 1280개의 스레드가 생성되어 돌아가기 때문에 잘못된 계산이 나올수 있으므로 예외 처리를 해줘야 한다. 하지만, add() 커널 함수에서 이미 예외 처리를 해놨었다.

그림 - 예외 처리가 된 부분

tid가 위에서 계산되었을 때 N보다 클 때에는 if문에서 걸러지기 때문에 올바르게 작동하게 된다.

하지만 이렇게 구현하게 되면, 하드웨어 제약사항으로 인한 블록의 개수(65,535) * 스레드의 개수(512) 개를 초과하게 되면, 병렬 처리가 정상적으로 처리가 되지 않게 된다. 그러므로 더 수정을 해야 한다.

그림 - 더 업그레이드 된 add() 커널 함수

위와 같이 각 스레드가 현재 인덱스에서의 작업을 끝내면 그리드 내에서 작동 중인 스레드의 총 수만큼 각 인덱스를 증가시킬 필요가 있다. 즉, while() 문에서 tid 값을 위와 같이 증가 시키면 각 스레드는 유니크한 인덱스를 구할 수 있을 것이다.

그림 - 고정된 개수의 블록과 스레드를 실행하는 add() 커널 함수

위와 같이 블록의 개수(128) * 스레드의 개수(128) = 16,384개 이지만, 3만개든 10만개든 덧셈이 동시에 가능해진다.