Synchronization

Race Condition 경쟁 상태

병렬 프로그램에서 Race condition (경쟁 상태)이 흔히 발생하는데 이는 여러 개의 스레드가 하나의 공유 메모리를 동시에 읽고 쓸 때 발생한다. 덧셈과 뺄셈같이 한 번에 이루어지는 계산은 아토믹 함수를 이용하고 더 넓은 범위의 코드는 크리티컬 섹션을 사용하여 경쟁 상태를 해결한다. 이 두 방법은 손쉬운 방법이지만 불필요하게 많이 사용하게 되면 오히려 프로그램의 효율이 떨어진다. 

 

//아래 코드 실행 시 10000*512개의 스레드가 실행되지만 Race condition(경쟁 상태) 발생으로 
//결과로 더 적은 수가 출력됨.
__global__ void ThreadCountDataRace(int* nThreadCount) {
	(*nThreadCount)++;
}

int main() {
	const int nBlocks = 10000; 
	const int nThreads = 512;

	int host_nThreadCount = 0;

	int* dev_nThreadCount;

	cudaMalloc((void**)&dev_nThreadCount, sizeof(int));

	cudaMemset(dev_nThreadCount, 0, sizeof(int));

	ThreadCountDataRace << <nBlocks, nThreads >> > (dev_nThreadCount);

	cudaMemcpy(&host_nThreadCount, dev_nThreadCount, sizeof(int),cudaMemcpyDeviceToHost);

	printf("실행 개수 : %d 개 \n", host_nThreadCount);

	cudaFree(dev_nThreadCount);

	return 0;
}

 

 

Atomic Functions 아토믹 함수

메모리 주소가 가리키는 값을 읽고 나서 연산을 수행한다. 같은 어드레스 결과를 저장하고 연산 완료 후 값을 반환한다. 

 

 

Thread Synchronization 스레드 동기화 : syncthreads

블록 내의 스레드를 동기화하도록 __syncthreads()함수 사용. 하나의 멀티프로세서에서 블록 단위로 스레드가 동작하기 때문에 __syncthreads()함수의 유효범위는 블록 내 스레드에 한정된다. 코드 내에 __syncthreads() 함수가 있으면 블록 내의 모든 스레드가 이 부분에 도달할 때 까지 대기한다.  __syncthreads()함수는 블록 내 스레드 동기화가 필요한 지점에 삽입하여 사용한다. 

 

 

Global Synchronizaton 글로벌 동기화

글로벌 동기화는 블록 단위의 동기화를 유지하기 위한 것으로, 글로벌 동기화를 지원하는 API가 따로 있는 것은 아니다. 글로벌 동기화를 유지하려면 커널 함수를 분리하여 호출하는 방법을 사용하며 커널을 반복 호출해야 하므로 커널 호출을 위한 부하가 추가된다. 글로벌 동기화를 위한 구현은 하나의 커다란 커널을 동기화가 필요한 지점별로 커널을 작게 나누어 여러 번 호출하거나 커널의 반복 호출을 하는 방법으로 이루어진다. 작은 커널로 나누어 구현할 때 의존성이 있는 데이터 간의 순서를 잘 지켜야 한다. 

//글로벌 동기화의 구현 방법
//1.동기화를 위한 커널의 분할
kernel_step1<<<nblocks, nThreads>>>(a,b); //a를 연산하여 b의 결과를 얻음.
kernel_step2<<<nblocks, nThreads>>>(b,c); //b를 연산하여 c의 결과를 얻음.

//2.커널의 반복 호출
for ( int i=0; i< limit, i++) 
{
	kernel<<<nblocks, nThreads>>>(input,result);
    input = result; 
}