GPU barrier: 데이터 경계와 참여 경계를 나누어야 하는 이유
GPU에서 일부 thread가 데이터를 채운 뒤 모두가 그 데이터를 읽으려면, ‘대충 같은 속도로 진행할 것’이라는 기대가 필요해 보인다. GPU barrier는 그 기대를 명시적인 참여와 순서의 계약으로 바꾸는 장치다. 다만 기다리는 범위와 도착할 thread를 잘못 정하면 동기화가 실행을 멈추게 할 수 있다.
같은 동기화 지점에 누가 도착해야 하는가
CUDA의 __syncthreads()는 block 안의 아직 종료하지 않은 thread들이 같은 호출 지점에 도착하거나 종료할 때까지 기다리며, 참여자 사이의 메모리 순서를 제공한다. 조건문 안에서 사용할 때는 문서의 block 전체에 대한 균일한 참여 조건을 지켜야 한다. CUDA Programming Guide, Thread Block Synchronization Functions
block에 128개 thread가 살아 있고 그중 앞의 64개만 동기화 호출에 들어간다고 하자. 나머지 64개가 종료하지도 않고 그 호출에 도착하지 않은 채 다른 작업을 계속한다면 필요한 참여가 완성되지 않는다. 조건이 thread마다 다른 이런 사용은 계약을 위반하며 hang이나 의도하지 않은 동작을 낳을 수 있다. 실제로 항상 같은 시간에 멈춘다는 실측 주장은 아니다.
block 전체가 같은 조건을 평가하여 모두 호출하거나 모두 건너뛰는 경우와 구별해야 한다. 또한 이미 종료한 thread를 무조건 살아 있는 참여자로 세어 설명하면 현재 문서의 의미와 어긋난다. ‘전체’라는 단어를 쓸 때에는 호출 지점에서의 종료 상태와 제어 흐름까지 확인한다.
데이터 경계와 참여 경계를 따로 정한다
가상의 단일 block은 128개 thread, 입력은 100개의 정수라고 하자. 목표는 입력의 순서를 뒤집는 것이다. shared 배열 s는 128개 원소를 가지고, 각 thread는 자신의 한 칸에만 쓴다.
먼저 thread 번호 t가 100보다 작으면 s[t]에 입력[t]를 쓰고, 나머지 28개는 자신의 칸에 0을 쓴다. 이 단계 이후 128개 thread 모두 같은 __syncthreads()를 호출한다. 그런 다음 t<100인 thread만 출력[t]에 s[99−t]를 쓴다. 출력 0은 입력 99, 출력 1은 입력 98, 출력 99는 입력 0을 받는다.
입력을 읽고 출력을 쓰는 조건은 유효 데이터 100개에 맞췄지만, barrier 참여는 block의 128개에 맞췄다. 배열 밖을 읽지 않으려는 조건을 동기화 호출 전체의 조건으로 잘못 사용하지 않는 것이 핵심이다. 이 예는 단일 block과 0≤n≤128의 구조를 지정한 추적이며 임의 크기의 다중 block reverse 구현은 아니다.
도착했다는 사실과 새로운 값이 보인다는 사실
메모리를 공유하는 협력에서는 실행 도착만이 아니라 읽기·쓰기 순서도 필요하다. 단순히 한 thread가 flag를 바꾸는 방식으로 barrier를 흉내 내면 flag의 원자성과 메모리 순서, 관측 범위를 추가로 정의해야 한다. flag 하나가 shared 배열의 안전한 전달을 자동으로 보장하지 않는다.
warp 범위만의 협력이라면 그 범위에 맞는 동기화와 참여 mask를 검토할 수 있다. 그러나 warp 하나의 동기화로 block의 다른 warp가 쓴 값까지 준비됐다고 결론 내릴 수는 없다. 데이터 생산자와 소비자가 속한 범위에 맞추어 계약을 선택한다.
또 block barrier는 grid 전체의 barrier가 아니다. CUDA의 일반 kernel launch는 block 실행 순서를 보장하지 않는다. CUDA Programming Guide, Kernel Launch and Occupancy
자체적인 스케줄 예에서 동시에 상주할 block 자리가 두 개뿐인데, 먼저 들어온 두 block이 아직 배치되지 않은 세 번째 block의 flag를 기다린다고 하자. 앞의 두 개가 자리를 놓지 않으면 세 번째가 생산할 기회도 없다. 메모리가 잘 보이는 flag를 사용해도 진행 조건의 문제는 남는다. grid 전체 협력에는 지원되는 실행·동기화 방식이 필요하다.
GPU barrier를 넣기 전에는 어느 thread가 어떤 값을 생산하고, 누가 같은 지점에 도착하며, 그 범위가 실행 가능한지 확인한다. 데이터의 유효 범위와 동기화 참여 범위를 나누어 설계하면 경계 처리 조건이 협력의 계약을 깨뜨리는 일을 줄일 수 있다.