CUDA Syncronization

CUDA 의 동기화에는 Atomic, Thread, Global 동기화 이렇게 3가지가 있다.

1. Atomic

 - 여러 Thread 에서 공유 메모리 변수에 동시에 읽고 쓰는 것을 막는 과정
    (CUDA에서는 Thread가 동시에 실행되므로, 이런 상황은 굉장히 빈번히 일어남)
 - 사용시 성능감소는 감수해야함. 가능하다면, 개별적으로 데이터 사용후 마지막에 결과값을 집계하도록 사용하는 것이 좋음
 - Atomic 함수 : Parameter 로 값이 저장된 주소, 새로운 Value 값을 받아서, 계산된 결과를 그 주소에 저장하고 그 값을 Return 한다.
                      (atomicCAS의 경우에는 2개의 Value를 입력받는데 앞에껀 compare, 뒤에껀 val 이다.)

atomicAdd : old + val
atomicSub : old - val
atomicExch : swap(old, val)
atomicMin : Min(old, val)
atomicMax : Max(old, val)
atomicInc : (old >= val) ? 0 : old+1
atomicDec : (old == 0 || old > val) ? val : old - 1
atomicCAS : (old == compare) ? val : old
atomicAnd : old & val
atomicOr : old | val
atomicEor : old ^ val

2. Thread 동기화

 - 블록 내의 스레드들을 동기화하는 것이며 __syncthreads() 함수를 사용한다.
 - 코드 내에 __syncthreads() 함수가 있으면 블록 내의 모든 스레드가 이 부분에 도달할 때까지 대기하게 된다.
 - if 등 조건문안에 __syncthreads()가 있을 경우 스레드들 중 해당 조건 안으로 못 들어가는 것이 발생할 경우 Deadlock이 발생 할 수도 있다.

3. Global 동기화

 - 블록 단위의 동기화이며, 지원하는 API가 따로 있는 것이 아니라, 커널 함수를 분리하여 호출하는 방법으로 구현

 - cudaThreadSyncronize() 를 이용하여 Kernel 함수가 완료될때까지 기다리는 방법을 이용하기도 한다.