inflearn logo
강의

講義

知識共有

CUDAプログラミング (5) - C/C++/GPU並列コンピューティング - アトミック演算 atomic op

30-4 histogram shared - shared memoryを使う

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

159

tizm4233217

投稿した質問数 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

onemoresipofcoffee

안녕하세요.

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

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

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

감사합니다.

0

tizm4233217

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

6강 17:50

0

10

1

6강 11:37

0

13

1

08:30 이중포인터 질문

0

10

0

코딩살구클럽 가입 요청드립니다.

0

11

2

6강 5:15

0

24

1

5강 5:45

0

16

2

전체적으로 답을 먼저 알려주지 마세요..

0

20

2

코딩살구클럽 채점 안됨

0

13

2

4-O 채점 관련 질문

0

20

2

코딩살구클럽가입

0

27

2

코딩살구클럽 가입

0

40

2

업캐스팅 문제 예시 5

0

42

2

(기출) 2025년 2회 이론 파트

0

41

2

26년 1회실기

0

45

2

비쥬얼 스튜디오로 해도 되나요?

0

52

2

강의 자료

0

40

2

scanf_s 에 관해서 오류나옵니다.

0

58

3

27:15 break 출력

0

40

2

최신 엔비디아 CUDA 아키텍처에서의 결정적 변경 사항

0

26

1

자문자답- 맞는지 틀린지 확인부탁드립니다.

0

34

1

55강 파이썬에만있는 연산자들

0

39

2

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

0

620

1

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

0

373

1

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

0

435

1