Ten Minute Physics - 16 Simulation on the GPU

GPU: 화면의 각 픽셀의 색상을 계산하기 위해 하나의 프로그램(shader)을 여러 객체에 적용

Simulation: 하나의 프로그램(kernel)을 수 많은 입자(particle)나 제약 조건(constraint)에 적용

Example: PBD Velocity Update

• 커널 정의 : @wp.kernel
• 스레드 ID 사용 : wp.tid() → 현재 스레드가 담당하는 입자 번호를 리턴
• 커널 실행 : wp.launch() → 커널 이름, 입력 버퍼, 스레드의 수, 실행할 장치 지정
• 데이터 이동 : wp.copy() → GPU에서 CPU로 데이터를 복사해 렌더링에 사용 혹은 초기 설정값을 보낼 때 사용.

GPU 병렬 처리 시 발생 가능한 문제

Challenge 1 - 동시 추가 (Simultaneous Adds)

• 제약 조건을 처리할 때에는 여러 스레드가 동일한 입자의 위치 수정값을 누적하려고 해 문제 발생
• wp.atomic_add() 와 같은 원자적 연산(atomic operations)을 사용해 문제 해결
  → 동일한 위치에 쓰기를 하고 있는 다른 스레드를 기다리기에 실행 속도가 약간 늦춰질 수 있음.
  

Challenge 2 - 동시 읽기 및 추가 (Simultaneous Read and Add)

• 일반적인 solver(Gauss-Seidel 방식)는 제약 조건을 해결하면서 입자 위치를 즉시 업데이트함.
• 하지만, GPU에서는 수많은 스레드가 정해지지 않은 순서로 실행됨.
  → 스레드 실행 순서에 따라, 어떤 제약 조건이 업데이트 되기 전의 위치를 읽을 수도 있고, 업데이트된 후의 위치를 읽을 수 있기에 시뮬레이션 결과가 비결정적(non-deterministic)이 되고, 출렁이는(jittering) 현상이 발생할 수 있게 됨.
  • 해결 방안 : Jacobi solver 또는 Graph Coloring

Jacobi Solver

• 입자 위치 $(x_i)$를 즉시 수정하는 대신, 수정 버퍼 $(d_i)$를 따로 사용함.
• 제약 조건 해결이 끝난 후, 모든 입자에 대해 $x_i \leftarrow x_i + sd_i$ 수식을 사용해 한 번에 수정 값 적용
• 장점: 모든 스레드가 솔버를 실행하는 동안 동일한 입자 위치로 작업하기에 비결정성이 사라짐
• 단점: Gauss-Seidel 방식보다 수렴 속도가 느리고, 진동(Oscillation)이나 발산이 발생할 수 있음. 따라서 수정 값에 1보다 작은 스칼라 $s$를 곱해줘야 함.

Graph Coloring

• 제약 조건들을 독립적인 부분 집합으로 나누어서 동작
  → 같은 색상의 제약 조건들은 서로 공유하는 입자가 없어서(Independent), 동시에 실행해도 충돌(Race condition)이 발생하지 않음.
  
• 안정적이며 Jacobi Solver처럼 $s$값을 선택할 필요가 없음
• 최적의 해를 찾는 것은 NP-hard 문제이기에, 일반적으론 Greedy algorithm 방식으로 해답을 찾음.

Hybrid Solution

• Graph Coloring을 사용하더라도 일반적으로 여러 번의 패스가 필요함.

• 또한 제약 조건 크기 분포를 확인하면, 초반에 큰 세트가 나오고 그 뒤로 작은 세트들의 긴 꼬리가 이어짐

  → 앞에 큰 세트들을 Graph Coloring으로 처리하고, 뒤에 긴 꼬리에 해당하는 작은 세트들은 Jacobi iteration으로 처리하여 시뮬레이션을 최적화