[30-4] hist-shared.cu 에서 질문있습니다.

CUDA 프로그래밍 (5) - C/C++/GPU 병렬 컴퓨팅 - 아토믹 연산 atomic op

30-4 histogram shared - shared memory 사용

작성

"hist-shared.cu"의 다음 부분에서 질문있습니다.

if(threadIdx.x < HIST_SIZE) {
    atomicAdd(&(hist[threadIdx.x]), s_hist[threadIdx.x]);
}

threadIdx.x는 0~31로 모두 다르기 때문에 병렬 연산이긴 해도, global memory에 있는 hist 배열의 서로 다른부분으로 write가 이루어진다고 생각했습니다.

따라서 서로 겹치는 위치가 없어서 atomic한 연산이 필요없을거라 생각했는데, atomic 연산이 없으면 결과가 이상하게 나오더라구요..

혹시 왜 그런지 알 수 있을까요?

양질의 강의 제공해주셔서 감사합니다!

답변 2

지식공유자

안녕하세요.

thread block 1개를 기준으로 보면, 말씀하신 내용이 맞습니다.

다만, 현재 예제에서는 thread block 을 여러 개 사용하고, 이것들이 동시에 실행되고 있습니다.

즉, thread block 여러 개에서, theadIdx.x 는 0 ~ 31인 thread 들이, 각 thread block 에서 동시에 메모리 영역에 write 하려고 하므로, 결국 atomicAdd 가 아니면 답이 틀리게 됩니다.

감사합니다.

질문자

이해했습니다 감사합니다!

인프런 커뮤니티 질문&답변