inflearn logo
강의

강의

N
챌린지

챌린지

멘토링

멘토링

N
클립

클립

로드맵

로드맵

지식공유

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

30-4 histogram shared - shared memory 사용

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

145

정재헌

작성한 질문수 4

0

"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 연산이 없으면 결과가 이상하게 나오더라구요..

 

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

 

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

c c++ cuda gpu 병렬-처리

답변 2

1

드립커피+한모금더

안녕하세요.

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

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

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

감사합니다.

0

정재헌

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

Reticle이 안나옵니다.

0

10

1

진행 방법 질문드립니다!

0

30

2

Singleton 관련 질문입니다.

1

31

2

안녕하세요. 계속 프로젝트를 해야지 하다가 결제하고 환경 설정 중입니다.

0

15

1

Export template 안됨

1

28

2

2-I) 왜 이 문제가 그래프이론 카테고리에 있는지 잘 모르겠습니다.

0

55

2

AddGameFrameworkComponentReceiver vs AddExtensionHandler

0

24

0

scanf("%d\n") 의미

0

20

1

필기자료 사라졌나요?(실기 일주일만에 안돼서 재도전-_-)

0

37

2

2주차 개념#12 트리 순회

0

25

2

백준 사이트 서비스 종료

0

112

3

26년 1회 실기 해설 강의

0

51

2

프로젝트 질문 문의

0

46

1

주소 연산자(&) 간접 지정자(*) 반대 개념

0

33

1

53번 4-1 자료 오류 있는 것 같습니다.

0

68

2

7번문제

0

58

2

C언어 변형문제 9번문제 Pdf 수정요청

0

45

2

메서드 오버드라드

0

45

2

실수

0

45

1

백준사이트가 종료된다고 합니다.

0

284

2

공부 우선순위 우선강의 알려주세요

0

85

1

선생님 gpu->cpu 속도 개선에 대해서 질문드려요

0

605

1

32-2의 warp의 early terminate에 대해 질문드립니다.

0

360

1

28-4 shuffle shared - shared memory 강의에서 질문이 있습니다.

0

422

1